Глобальная запись в память считается атомарной или нет в CUDA?

Учитывая следующий код ядра CUDA:

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

Глобальная память пишет вglobalStorage атомный ?, напримернет условий гонки, при которых параллельные потоки ядра записывают в байты одной и той же переменной, хранящейся в globalStorage, что могло бы испортить результаты (например, записи с частичной записью)?

Обратите внимание, что здесь я не говорю об элементарных операциях, таких как add / sub / bit-bit и т. Д., Просто о глобальной записи.

Отредактировано: переписан пример кода, чтобы избежать путаницы.