CUDA支持指针混叠吗?

Question

我问这个问题的原因是因为我的代码中有一些奇怪的错误，我怀疑这可能是一些混叠问题:

__shared__ float x[32];
__shared__ unsigned int xsum[32];
int idx=threadIdx.x;
unsigned char * xchar=(unsigned char *)x;
//...do something
 if (threadIdx.x<32)
 {
    xchar[4*idx]&=somestring[0];
    xchar[4*idx+1]&=somestring[1];
    xchar[4*idx+2]&=somestring[2];
    xchar[4*idx+3]&=somestring[3];
    xsum[idx]+=*((unsigned int *)(x+idx));//<-Looks like the compiler sometimes fail to recongize this as the aliasing of xchar;
 };

Answer 1

编译器只需要支持兼容类型之间的混叠。由于char和float不兼容，编译器可以自由地假设指针没有别名。

如果您想对float进行位操作，首先转换(通过__float_as_int())为无符号整数，然后对其进行操作，最后转换回float(使用__int_as_float())。

Answer 2

我想你这里有一个竞态条件。但是我不知道什么是somestring。如果所有线程都是相同的，你可以这样做:

__shared__ float x[32];
unsigned char * xchar=(unsigned char *)x;
//...do something
if(threadIdx.x<4) {
     xchar[threadIdx.x]&=somestring[threadIdx.x];
}
__syncthreads();
unsigned int xsum+=*((unsigned int *)x);

这意味着每个线程共享同一个数组，因此，所有线程之间的xsum是相同的。如果你想让每个线程都有自己的数组，你必须分配一个32*number_of_threads_in_block数组并使用偏移量。

PS:以上代码只适用于1D块。在2D或3D中，你必须计算自己的threadID，并确保只有4个线程执行代码。