多個變量在CUDA中同步

Question

我的程序有很多4字節字符串，比如「aaaa」「bbbb」「cccc」...我需要收集通過crc檢查的特定字符串。

因爲字符串可以通過crc檢查的可能性很小，所以我不想使用非常大的緩衝區來保存所有結果。我喜歡一個接一個的結果，就像輸入一樣。例如，如果輸入是「aaaabbbbcccc」和「BBBB」未通過CRC校驗，則輸出字符串應該是「aaaacccc」和output_count應該爲2

代碼看起來像：

__device__ 
bool is_crc_correct(char* str, int len) { 
    return true; // for simplicity, just return 'true'; 
} 

// arguments: 
// input: a sequence of 4-bytes-string, eg: aaaabbbbccccdddd.... 
__global__ 
void func(char* input, int* output, int* output_count) { 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x; 

    if(is_crc_correct(input + 4*index)) { 
     // copy the string 
     memcpy(output + (*output_count)*4, 
       input + 4*index, 
       4); 
     // increase the counter 
     (*output_count)++; 
    } 
} 

顯然內存拷貝不是線程安全的，我知道atomicAdd函數可以用於++操作，但是如何使output和output_count線程安全？

Answer 1

你在找什麼是一個無鎖的線性分配器。這樣做的通常方法是通過原子級增加的累加器來索引緩衝區。例如，你的情況，下面應該工作：

__device__ 
char* allocate(char* buffer, int* elements) { 
    // Here, the size of the allocated segment is always 4. 
    // In a more general use case you would atomicAdd the requested size. 
    return buffer + atomicInc(elements) * 4; 
} 

然後可以使用這樣的：

__global__ 
void func(char* input, int* output, int* output_count) { 
    unsigned int index = blockDim.x*blockIdx.x + threadIdx.x; 

    if(is_crc_correct(input + 4*index)) { 
     // Reserve the output buffer. 
     char* dst = allocate(output, output_count); 
     memcpy(dst, input + 4 * index, 4); 
    } 
} 

雖然這是完全線程安全的，這是不保證保護輸入順序。例如，「ccccaaaa」將是一個有效的輸出。

---

如滴在他們的評論提到，你正在嘗試做的是一個有效的數據流壓縮（和推力已經可能已經提供了你所需要的）。

我上面發佈的代碼可以進一步優化，首先將輸出字符串以變形聚合，而不是直接分配到全局緩衝區。這將減少全球原子爭奪並可能導致更好的表現。有關如何執行此操作的解釋，請您閱讀以下文章：CUDA Pro Tip: Optimized Filtering with Warp-Aggregated Atomics。

Answer 2

我可能最終會暗示這一點，但如何在內核中動態分配內存？例如：CUDA allocate memory in __device__ function

然後，您可以將共享內存數組傳遞給每個內核，並且在內核運行後，數組的每個元素都將指向一塊動態分配的內存，或者空值。因此，在線程塊運行後，您將在單個線程上運行最終的清理內核來構建最終的字符串。