Поддерживает ли 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

Я думаю, что у вас есть условие гонки здесь. Но я не знаю, что такое 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 потока выполняют код.

Answer 2

Компилятор должен выполнять только псевдонимы между совместимыми типами. Поскольку char и float несовместимы, компилятор может предположить, что указатели никогда не являются псевдонимами.

Если вы хотите выполнять побитовые операции над float, сначала конвертируйте (через __float_as_int()) в unsigned integer, затем оперируйте его и, наконец, преобразуйте обратно в float (используя __int_as_float()).