CUDA inline PTX ld.shared 遇到 cudaErrorIllegalAddress 错误答案

【问题标题】：CUDA inline PTX ld.shared runs into cudaErrorIllegalAddress errorCUDA inline PTX ld.shared 遇到 cudaErrorIllegalAddress 错误
【发布时间】：2021-12-24 10:03:36
【问题描述】：

我正在使用内联 PTX ld.shared 从共享内存中加载数据：

__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];  //declare a buffer in shared memory
float Csub = 0;

As[TY][TX] = A[a + wA * TY + TX];             //load data from global memory to shared memory
__syncthreads();
float t;
asm("ld.shared.f32 %0, [%1];" :"=f"(t) : "r"((int)&As[TY][k]));  //load data from shared memory into t
Csub += t;
__syncthreads();

但遇到错误：

C:/ProgramData/NVIDIA Corporation/CUDA Samples/v11.2/0_Simple/matrixMul_mine/matrixMul.cu:196 code=700(cudaErrorIllegalAddress) "cudaStreamSynchronize(stream)" 出现 CUDA 错误

我转储 SASS 并发现 LDS 发生的时间甚至早于 LDG 和 2 bar.sync。好像编译器丢失了数据依赖的轨迹；

所以我的问题是：

我的内联 PTX 中是否有任何错误导致 cudaErrorIllegalAddress？
内联 PTX 会干扰数据依赖性吗？

【问题讨论】：

docs.nvidia.com/cuda/inline-ptx-assembly/…
也许除了 talonmies 答案之外，您还必须将 As 声明为 volatile，非法地址可能与订单无关 - 但我们看不到您的参数 a、wA、TY、TK、k 的位置来自，全局数组有多大，或者你的内核被调用的块和网格大小。
刚刚得知地址需要像弯刀中的link一样用"cvta"转换。正如塞巴斯蒂安所说，应该在“asm”之后添加“volatile”。

标签： cuda inline shared-memory ptx

【解决方案1】：

一尘是对的。

有两种寻址方式：ld。或 ld.statespace。

如果 ld.只是，地址应该是通用地址。据我了解（有限），通用地址是 CUDA-C 指针值，如代码中的 &As[TY][k]。

如果是ld.statespace，地址应该是状态空间中的地址。

我认为如果您使用 ld.f32 而不是 ld.shared.f32，您的代码应该没问题。顺便说一句，我认为您不能使用 32 位数据宽度的通用地址，这可能会将通用地址截断为错误的值。

或者您可以将通用地址转换为共享空间地址。这是弯刀的转换代码：

      ".reg .u32 smem_ptr32;\n\t"
      ".reg .u64 smem_ptr64; cvta.to.shared.u64 smem_ptr64, %1; cvt.u32.u64 smem_ptr32, smem_ptr64; \n\t"

然后使用 smem_ptr32 代替 [%1] “ld.shared.f32 %0, [smem_ptr32];” 正如 PTX ISA 所说，此地址可以是 32 位或 64 位。我认为没有必要将 64 位 ptr 转换为 32 位 ptr。使用 smem_ptr64 应该没问题。

这是共享内存地址的样子：

CUDA-C 指针（通用）：1526743433216 + 1024
smem_ptr64(共享空间):0 + 1024

【讨论】：