A gravação de memória global é considerada atômica ou não em CUDA?

Considerando o seguinte código de kernel CUDA:

int idx = blockIdx.x*blockDim.x+threadIdx.x;
int gidx = idx%1000;
globalStorage[gidx] = somefunction(idx);

A memória global escreve paraglobalStorage atómico ?, p.não há condições de corrida tais que os threads de kernel simultâneos escrevam para os bytes da mesma variável armazenada em globalStorage, o que poderia atrapalhar os resultados (por exemplo, gravações pariais)?

Note que eu não estou falando sobre operações atômicas como add / sub / bit-wise etc aqui, apenas escrita global direta.

Editado: reescreveu o código de exemplo para evitar confusão.