CUDA:结构体的共享数据成员和对该结构体的引用成员具有不同的地址和值
CUDA: shared data member of struct and member of reference to that struct have different addresses, values
好的,问题来了:
使用CUDA 1.1计算GPU,我试图为每个线程维护一组(可能变化的数量,这里固定为4)索引,作为结构变量的成员保留的引用
我的问题是,在访问成员数组时,获取对结构的引用会导致不正确的结果:我用0初始化成员数组值,当我使用原始结构变量读取数组值时,我得到了正确的值(0),但当我使用对结构var的引用读取时,我会得到垃圾(-8193)。即使使用class
而不是struct
,也会发生这种情况。
为什么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内存检查器捕获全局内存的越界存储。
Memory Checker detected 1 access violations.
error = access violation on store (global memory)
blockIdx = {0,0,0}
threadIdx = {0,0,0}
address = 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编译器应该将CCD_ 11转换为通用指针。
相关文章:
- 根据用户回答声明"Players"。用户选择玩家数量。播放器是结构体
- 结构体 S { int align; } 之间的区别;(struct 关键字后的名称)和 struct { int al
- C++ - 如何在结构向量中找到结构体一个成员的最大值?
- 包含 std::list 的结构体的 C++ 初始化
- 结构体和类的不同大小(),彼此具有相同的字段类型
- 如何使用结构体的向量数组?
- 如何使用结构体在C++中更改这些代码?
- 无法在 Mosquitto MQTT Broker 插件上访问结构体 mosquitto 的元素
- 我应该如何在C++中使用结构体解决输入失败的问题?
- Qsort() 比较结构体整数的总和
- 如何使用迭代器指向结构体c++的向量
- 在C++中使用链表的堆栈实现中,访问结构体headNode成员count和top会导致运行时错误
- 如何获取结构体成员的地址
- 创建结构体向量,表达式:向量下标超出范围
- boost::任何带有结构体和无符号整数
- 如何在构造函数中初始化结构体的动态数组?
- 只写结构体的某些字段
- CUDA:结构体的共享数据成员和对该结构体的引用成员具有不同的地址和值
- 如果我使用一个结构体来携带argc和argv,我如何将地址argv赋值给结构体中的变量?
- 如何在特定地址声明结构体