c++ - CUDA:结构的共享数据成员和对该结构的引用成员具有不同的地址、值

Question

好的，问题来了: 使用 CUDA 1.1 计算 gpu，我试图为每个线程维护一组(可能数量不同，这里固定为 4)索引， 我作为 struct var 的成员保留的引用。 我的问题是获取对结构的引用然后在访问成员数组时导致不正确的结果:我用 0 初始化成员数组值，当我使用原始结构 var 读取数组 vals 时，我得到正确的值 (0) ，但是当我使用对 struct var 的引用读取它时，我得到了垃圾 (-8193)。 即使使用类而不是结构，也会发生这种情况。

为什么 tmp 低于 != 0 ??

c++ 不是我的主要语言，所以这可能是一个概念问题，也可能是在 cuda 中工作的一个怪癖

struct DataIdx {
    int numFeats;
    int* featIdx;
};
extern __shared__ int sharedData[];

__global__  void myFn(){
    int tidx = blockIdx.x * blockDim.x + threadIdx.x;

    DataIdx myIdx;  //instantiate the struct var in the context of the current thread
    myIdx.numFeats = 4;
    size_t idxArraySize = sizeof(int)*4;
    //get a reference to my array for this thread. Parallel Nsight debugger shows myIdx.featIdx address = 0x0000000000000000e0
    myIdx.featIdx = (int*)(&sharedData[tidx*idxArraySize]);  

    myIdx.featIdx[0] = 0x0;  //set first value to 0 
    int tmp = myIdx.featIdx[0];  // tmp is correctly eq to 0 in Nsight debugger -- As Expected!!
    tmp = 2*tmp;    antIdx.featIdx[0] = tmp; //ensure compiler doesn't elide out tmp

    DataIdx *tmpIdx = &myIdx;  //create a reference to my struct var
    tmp = tmpIdx.featIdx[0];   // expected 0, but tmp = -8193 in debugger !! why?  debugger shows address of tmpIdx.featIdx = __devicea__ address=8
    tmpIdx.featIdx[0] = 0x0;
    tmp = tmpIdx.featIdx[0]; // tmp = -1; cant even read what we just set

    //forcing the same reference as myIdx.featIdx, still gives a problem! debugger shows address of tmpIdx.featIdx = __devicea__ address=8
    tmpIdx->featIdx =  (int*)(&sharedData[tidx*idxArraySize]); 
    tmp = tmpIdx.featIdx[0]; //tmp = -8193!! why != 0?

    DataIdx tmpIdxAlias = myIdx;
    tmp = tmpIdx.featIdx[0]; //aliasing the original var gives correct results, tmp=0


     myIdx.featIdx[0] = 0x0;
     mySubfn(&myIdx); //this is a problem because it happens when passing the struct by reference to subfns
     mySubfn2(myIdx);
}
__device__ mySubfn(struct DataIdx *myIdx){
  int tmp = myIdx->featIdx[0]; //tmp == -8193!! should be 0
}
__device__ mySubfn2(struct DataIdx &myIdx){
  int tmp = myIdx.featIdx[0]; //tmp == -8193!! should be 0
}

Answer 1

我必须修改你的代码才能编译。在行中

tmpIdx->featIdx[0] = 0x0

编译器无法理解指向共享内存的指针。它不是对共享内存 (R2G) 进行存储，而是对超出范围的全局地址 0x10 进行存储。

    DataIdx *tmpIdx = &myIdx;
0x000024c8  MOV32 R2, R31;  
0x000024cc  MOV32 R2, R2;  
    tmp = tmpIdx->featIdx[0];
    tmpIdx->featIdx[0] = 0x0;
0x000024d0  MOV32 R3, R31;  
0x000024d4  MOV32 R2, R2;  
0x000024d8  IADD32I R4, R2, 0x4;  
0x000024e0  R2A A1, R4;  
0x000024e8  LLD.U32 R4, local [A1+0x0];  
0x000024f0  IADD R4, R4, R31;  
0x000024f8  SHL R4, R4, R31;  
0x00002500  IADD R4, R4, R31;  
0x00002508  GST.U32 global14 [R4], R3;   // <<== GLOBAL STORE vs. R2G (register to global register file)
    tmp = tmpIdx->featIdx[0];

Nsight CUDA 内存检查器捕捉到全局内存的越界存储。

Memory Checker 检测到 1 次访问冲突。 错误 = 存储访问冲突(全局内存) blockIdx = {0,0,0} threadIdx = {0,0,0} 地址 = 0x00000010 访问大小 = 0

如果您为 compute_10,sm_10(实际上是 <= 1.3)编译，您应该会在编译器无法确定访问共享内存的每一行中看到以下警告:

kernel.cu(46): warning : Cannot tell what pointer points to, assuming global memory space

如果您在启动后添加 cudaDeviceSynchronize，您应该会看到由越界内存访问引起的错误代码 cudaErrorUnknown。

__shared__ 是一个可变内存限定符而不是类型限定符，所以我知道你会如何告诉编译器 featIdx 将始终指向共享内存。在 CC >= 2.0 上，编译器应将 (int*)(&sharedData[tidx*idxArraySize]) 转换为通用指针。