Issues with atomicAdd when counting chars


I’m trying to count how many times the char exists in a array.

The kernel that I’m using is this one.

[codebox]global void charCounter(char *string, int *charCounter, int N)


int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < N) 

	atomicAdd(&charCounter[(int)string[idx]], 1);



The result with the serial code is that all the characters got the same number. But the parallel results are different they are near the correct ones.

Mind that string being of char type will result in indices between -128 and +127 for your charCounter array.
Wouldn’t it be safer to use unsigned char instead?

Yes. It would be safer. It worked! :wave: