Константа для цикла дает неверный результат

Question

Я собираюсь написать ядро ​​гистограммы для устройств CUDA. Он основан на NVIDIA's paper.

Идея состоит в том, что каждая нить вычисляет частичную гистограмму определенной части (в моем случае тома) и записывает ее в блок разделяемой памяти. Тем не менее, я столкнулся странные проблемы с алгоритмом и урезанная ядра до значительных частей:

__global__ void calcHist64() 
{ 
    extern __shared__ unsigned char partialHistograms[]; 

    //a unique sequential thread id within this block, used to determine the memory in which to write the partial histogram 
    unsigned int seqTid = threadIdx.x + threadIdx.y * blockDim.x;  
#pragma unroll 
    for(int i = 0; i < 255; ++i) 
    { 
     //increment the thread's partial histogram value 
     partialHistograms[seqTid]++; 
    } 
    //each partial histogram should now be 255 
    //Output the value for every thread in a certain block 
    if(blockIdx.x == 0 && blockIdx.y == 31) 
     printf("Partial[%i][%i]: %i\n", threadIdx.x, threadIdx.y, partialHistograms[partialHistRoot]); 
} 

Ядро вызывается через:

int sharedMemory = 4096; 
dim blocks(32, 32, 1); 
dim3 threadsPerBlock(8,8,1); 
calcHist64<<<blocks, threadsPerBlock, sharedMemory>>>(); 

Я ожидаю, что каждый частичный гистограмму иметь значение 255. Однако это справедливо только для первых блоков (низкий blockIdx.x/blockIdx.y). Значения для других блоков сильно различаются. Последние блоки (blockIdx.y == 31) имеют значения 239 или 240.

Я не могу объяснить это поведение. Это константа для цикла, которая выполняется ровно 255 раз. Каждый поток обращается к различным частям разделяемой памяти, поэтому не должно быть никаких условий гонки.

Может ли кто-нибудь объяснить это поведение?

Answer 1

Вы не инициализируете общую память. Он автоматически не инициализируется нулем для вас.

После этой строки кода:

unsigned int seqTid = threadIdx.x + threadIdx.y * blockDim.x; 

Добавить это:

partialHistograms[seqTid] = 0; 

Также в вашем коде вы не определяете partialHistRoot. Я предполагаю, что partialHistRoot == seqTid. Если это не так, у вас есть условие гонки и ваше заявление

Каждый поток обращается к различным частям общей памяти, поэтому не должно быть условий гонки.

Неправильная версия.