Issue with running kernel on blackwell GPUs (RTX 5090)

[huoshenlaile]

Hello, I tried to run the Marlin kernel on a Blackwell GPU, specifically RTX 5090, with CUDA 12.8. and got RuntimeError: CUDA error: an illegal instruction was encountered. Is there a way to fix this? Thanks in advance.

[DrXuQian]

Are you able to solve this issue?

[huoshenlaile]

Are you able to solve this issue?

Yes, currently my fix is to just remove the cache policy for async copy of B matrix:
Replace the cp_async4_stream with the code below.

__device__ inline void cp_async4_stream(void* smem_ptr, const void* glob_ptr) {
  const int BYTES = 16;
  uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
  asm volatile(
    "{\n"
    "   cp.async.cg.shared.global [%0], [%1], %2;\n"
    "}\n" :: "r"(smem), "l"(glob_ptr), "n"(BYTES)
  );
}