塊級原子寫

Question

是否可以在塊級上執行原子寫入？
作爲一個例子考慮以下因素：

__global__ kernel (int atomic) 
{ 
    atomic+=blockid.x; //should be atomic for each block 
} 

Answer 1

雖然不清楚你的意思塊/塊級的東西，這聽起來像你只需要一個原子添加。 這些都在#include <asm/atomic.h>代碼內核發現將

__global__ kernel (int atomic) 
{ 
    atomic_add(blockid.x,&atomic); 
} 

atomic將不得不atomic_t類型和blockid.x一個int。

Answer 2

您可以在CUDA中執行一些原子操作。請參見CUDA編程指南中的附錄B.11原子函數。即：

__global__ void kernel (int *result) 
{ 
    atomicAdd(result, blockIdx.x); // 
} 

您還可以交換的變量

__global__ void kernel (int *result) 
{ 
    atomicExch(result, blockIdx.x); // 
} 

兩個例子在全局內存操作價值。

Atomic functions operating on shared memory and atomic functions operating on 64-bit words are only available for devices of compute capability 1.2 and above.

問候。

Answer 3

您可以對共享內存執行原子操作，但不是您在代碼片段中嘗試這麼做的方式：您的內核的int參數是特定於線程的變量;即使所有線程在啓動時都獲得了相同的值，但它們不會將其存儲在共享內存中 - 並且以原子方式對其進行操作是沒有意義的。

如果您已將某個int *傳遞給某個緩衝區 - 這可能是全局內存中的緩衝區。您可以對全局內存中的數據執行設備範圍內的原子操作，如@ pQB的answer中所述。但是你問到塊級原子操作......對全局數據沒有多大意義。但是，如果其中一個線程寫入某個全局地址，則全部可以掛起__threadfence_block()，直到該寫入的結果對該塊中的所有其他線程都可見。

CUDA也支持正確的塊級原子，但是 - 在共享內存上。請閱讀如何在this Parallel4All blog entry或相關的sectionCUDA Programming Guide中使用共享內存。

如果你有一些__shared__ int x，你確實可以對它執行一個塊級的原子操作，其語法與全局原子相同：atomicAdd(&x, 7)會自動將123添加到x的值。但是 - 請記住塊中的所有線程都會執行相同的操作，並且您當然不希望一次嘗試多達1024個原子寫入。通常情況下，你會有類似的東西

__shared__ some_buffer[NumFoosPerBar]; 

// ... 

if (check_condition()) { 
    int foo_index = get_thread_foo_index_for(threadIdx.x); 
    atomicAdd(&some_buffer[foo_index], 7); 
} 

其中可能有多個線程寫入相同的位置，但不一定。當你做期望多次寫入 - 不要使用原子，而是執行某些減少​​寫入的值。