Can we call CUTLASS GEMM operations within cuda kernel, but not at the host code?

[jchan192]

I am looking for a cuda GEMM operation being called within the kernel, but not at the host code. Several discussion ,llke this recommend us to look at CUTLASS. But I cannot see any documentation/example allows calling CUTLASS operation within a kernel. Do you mind to provide more info whether doing the above is possible using CUTLASS. Thank you for your time.

For example

__global__ void Func(A,B){

//  doing thing on data A and B... 

// perform GEMM using CULASSS? 

return C
}

[jackkosaian]

One route to achieving this is by calling into the cutlass::Kernel function (see here) for CUTLASS 2.x kernels or the cutlass::device_kernel function (see here) for CUTLASS 3.x kernels. These are used by the __global__ functions for launching kernels in CUTLASS.

Each of these functions takes as a template parameter a kernel::-scoped operator (e.g., cutlass::gemm::kernel::GemmUniversal here). To use one of these functions, you'll need to construct the Params input argument to the kernel::-scoped operator's operator() method (e.g., here for GemmUniversal). The constructed Params will be passed into the Kernel or device_kernel function.

Since you'll be bypassing many of the host-side utilities provided by CUTLASS's device::-scoped GEMMs, you'll need to manually do the following:

• You will need to launch a grid that is the same size as that which would be launched by the device::-scoped GEMM (e.g., that constructed here).
• You will need to make sure that the amount of dynamic shared memory available can support the amount required by your CUTLASS kernel. See the code here for how to achieve this.
• You will need to allocate and initialize any workspace needed by the CUTLASS kernel. You can see examples of how workspace size is determined in CUTLASS's host-side utilities here.

I hope this helps!