[QST] [CuTeDSL] DSLCudaRuntimeError: Unknown CUDA error

[HanGuo97]

What is your question?
The following code (which is essentially an empty kernel) will trigger DSLCudaRuntimeError: Unknown CUDA error.

import cutlass
import cutlass.cute as cute

@cute.kernel
def kernel():
    layout = cute.make_layout((16, 16))
    smem = cutlass.utils.SmemAllocator()
    smem.allocate_tensor(element_type=cute.Float32, layout=layout)
    


@cute.jit
def device():
    kernel().launch(grid=(1,1,1), block=(1,1,1))

device()

The above code will lead to

DSLCudaRuntimeError: DSLCudaRuntimeError: Unknown CUDA error
Error Code: 201

🔍 Additional Context: 
- Error name: CUDA_ERROR_INVALID_CONTEXT
- CUDA_TOOLKIT_PATH: not set
- Target SM ARCH: not set

📊 GPU Information:
- CUDA devices available: 8 (current: 0)
- Architecture: Ampere (sm_80)
- Compatible SM archs: sm_80

Compatibility Check:
❌ Error: Target SM ARCH unknown is not compatible
💡 Please use one of SM ARCHs: sm_80

Interestingly, if I did ran another working kernel, the same code will work again.

[fengxie]

It's because cuda context is not initialized. This should fix the problem. ( Will update error message to provide actionable information ). @brandon-yujie-sun

cutlass.cuda.initialize_cuda_context()
device()

Another thing, it's recommended to set shared memory size. ( We are working on an improvement to automatically set shared memory. Maybe one or two releases away. Once it's done, smem=xxx will not be requried.

@cute.jit
def device():
    kernel().launch(grid=(1,1,1), block=(1,1,1), smem=16 * 16 * 4)

[HanGuo97]

Ah got it, thank you! When do we need to use initialize_cuda_context? Somehow a few kernel programs work without it so I didn't realize this is needed

[fengxie]

Ah got it, thank you! When do we need to use initialize_cuda_context? Somehow a few kernel programs work without it so I didn't realize this is needed

If you are working under torch, you don't need it in most cases as torch will initialize it. Or any CUDA API is called before this kernel launch. In your example, there is no CUDA API called so no one initialize CUDA context.

E.g. if you allocate a buffer on GPU with CUDA API, it will be created. But if you are using cuda-python, most likely, you need to call cuda.cuInit(0) before any other cuda API can be called.

An analogy in C++

• CUDA runtime API usually will initialize/manage context automatically with the first API call
• CUDA driver API don't create context so people usually call cudaFree(0) to do this trick.

CuTe DSL is design in a way to do minimum staff to avoid interference with context managed by framework or application. It doesn't manage context.

[HanGuo97]

This is very clear now, thank you! Most prior kernels I used involved PyTorch tensors, hence I didn't realize this.