CUDA:struct的共享数据成员和该struct的引用成员具有不同的地址,值
Posted
tags:
篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了CUDA:struct的共享数据成员和该struct的引用成员具有不同的地址,值相关的知识,希望对你有一定的参考价值。
接下来是问题:使用CUDA 1.1计算gpu,我试图为每个线程维护一组(可能不同数量,这里固定为4)索引,这是我作为struct var成员保留的引用。我的问题是获取对结构的引用然后在访问成员数组时导致不正确的结果:我初始化成员数组值为0,当我使用原始struct var读取数组val时,我得到正确的值(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
}
我不得不修改你的代码来编译。在线
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内存检查器将超出范围的存储捕获到全局内存。
内存检查器检测到1次访问冲突。 error =存储上的访问冲突(全局内存)blockIdx = {0,0,0} threadIdx = {0,0,0}地址= 0x00000010 accessSize = 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])转换为通用指针。
以上是关于CUDA:struct的共享数据成员和该struct的引用成员具有不同的地址,值的主要内容,如果未能解决你的问题,请参考以下文章