正确进行多个CUDA块同步的方法
The way to properly do multiple CUDA block synchronization
我喜欢为多个块做CUDA同步。并不是每个 __syncthreads() 可以轻松处理的块。
我看到有关于这个主题的现有讨论，例如 cuda block synchronization, and I like the simple solution brought up by @johan, ，本质上它使用 64 位计数器来跟踪同步块。
然而，我写了下面的代码试图完成类似的工作但是遇到了一个问题。这里我使用了术语 environment 以便此环境中的块 wkNumberEnvs 应该同步。它有一个柜台。我使用 atomicAdd() 来计算自己已经同步了多少个块，一旦同步块的数量 == wkBlocksPerEnv，我就知道所有块都已完成同步并且可以自由使用。但是，它有一个奇怪的结果，我不确定为什么。
问题出在这个while循环中。由于所有块的第一个线程都在执行 atomicAdd，因此有一个 while 循环来检查直到条件满足。但是我发现有些block会陷入死循环，不知道为什么最终不能满足条件？如果我在 *** I can print here 1 或 *** I can print here 2 中打印一些消息，就没有无限循环，一切都很完美。我没有看到明显的东西。
const int wkBlocksPerEnv = 2;
__device__ int env_sync_block_count[wkNumberEnvs];
__device__ void syncthreads_for_env(){
// sync threads for each block so all threads in this block finished the previous tasks
__syncthreads();
// sync threads for wkBlocksPerEnv blocks for each environment
if(wkBlocksPerEnv > 1){
const int kThisEnvId = get_env_scope_block_id(blockIdx.x);
if (threadIdx.x == 0){
// incrementing env_sync_block_count by 1
atomicAdd(&env_sync_block_count[kThisEnvId], 1);
// *** I can print here 1
while(env_sync_block_count[kThisEnvId] != wkBlocksPerEnv){
// *** I can print here 2
}
// Do the next job ...
}
}
原子值将进入全局内存，但在 while-loop 中您直接读取它并且它必须来自不会在线程之间自动同步的缓存（cache-coherence 仅由显式同步处理像 threadfence）。线程获得自己的同步，但其他线程可能看不到它。
即使您使用 threadfence，同一个 warp 中的线程也会 dead-lock 如果它们是第一个在任何其他线程更新它之前检查该值的线程，那么它们将永远等待。但应该与支持独立线程调度的最新 GPU 一起工作。
您的代码存在两个潜在问题。缓存和块调度。
缓存可以防止您在 while 循环期间观察更新的值。
如果您等待尚未安排的块更新，则块安排可能会导致 dead-lock。由于 CUDA 不保证调度块的特定顺序，因此防止此 dead-lock 的唯一方法是限制网格中的块数，以便所有块可以同时 运行。
以下代码显示了如何在避免上述问题的同时同步多个块。我从 CUDA-sample conjugateGradientMultiDeviceCG https://github.com/NVIDIA/cuda-samples/blob/master/Samples/4_CUDA_Libraries/conjugateGradientMultiDeviceCG/conjugateGradientMultiDeviceCG.cu#L186
中给出的 multi-grid 同步改编了代码
在 pre-Volta 设备上，它使用易失性内存访问。 Volta 和后来使用 acquire/release 语义。
网格大小受查询设备属性限制。
#include <cassert>
#include <cstdio>
constexpr int wkBlocksPerEnv = 13;
__device__
int getEnv(int blockId){
return blockId / wkBlocksPerEnv;
}
__device__
int getRankInEnv(int blockId){
return blockId % wkBlocksPerEnv;
}
__device__
unsigned char load_arrived(unsigned char *arrived) {
#if __CUDA_ARCH__ < 700
return *(volatile unsigned char *)arrived;
#else
unsigned int result;
asm volatile("ld.acquire.gpu.global.u8 %0, [%1];"
: "=r"(result)
: "l"(arrived)
: "memory");
return result;
#endif
}
__device__
void store_arrived(unsigned char *arrived,
unsigned char val) {
#if __CUDA_ARCH__ < 700
*(volatile unsigned char *)arrived = val;
#else
unsigned int reg_val = val;
asm volatile(
"st.release.gpu.global.u8 [%1], %0;" ::"r"(reg_val) "l"(arrived)
: "memory");
// Avoids compiler warnings from unused variable val.
(void)(reg_val = reg_val);
#endif
}
#if 0
//wrong implementation which does not synchronize. to check that kernel assert does trigger without proper synchronization
__device__
void syncthreads_for_env(unsigned char* temp){
}
#else
//temp must have at least size sizeof(unsigned char) * total_number_of_blocks in grid
__device__
void syncthreads_for_env(unsigned char* temp){
__syncthreads();
const int env = getEnv(blockIdx.x);
const int blockInEnv = getRankInEnv(blockIdx.x);
unsigned char* const mytemp = temp + env * wkBlocksPerEnv;
if(threadIdx.x == 0){
if(blockInEnv == 0){
// Leader block waits for others to join and then releases them.
// Other blocks in env can arrive in any order, so the leader have to wait for
// all others.
for (int i = 0; i < wkBlocksPerEnv - 1; i++) {
while (load_arrived(&mytemp[i]) == 0)
;
}
for (int i = 0; i < wkBlocksPerEnv - 1; i++) {
store_arrived(&mytemp[i], 0);
}
__threadfence();
}else{
// Other blocks in env note their arrival and wait to be released.
store_arrived(&mytemp[blockInEnv - 1], 1);
while (load_arrived(&mytemp[blockInEnv - 1]) == 1)
;
}
}
__syncthreads();
}
#endif
__global__
void kernel(unsigned char* synctemp, int* array){
const int env = getEnv(blockIdx.x);
const int blockInEnv = getRankInEnv(blockIdx.x);
if(threadIdx.x == 0){
array[blockIdx.x] = 1;
}
syncthreads_for_env(synctemp);
if(threadIdx.x == 0){
int sum = 0;
for(int i = 0; i < wkBlocksPerEnv; i++){
sum += array[env * wkBlocksPerEnv + i];
}
assert(sum == wkBlocksPerEnv);
}
}
int main(){
const int smem = 0;
const int blocksize = 128;
int deviceId = 0;
int numSMs = 0;
int maxBlocksPerSM = 0;
cudaGetDevice(&deviceId);
cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceId);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&maxBlocksPerSM,
kernel,
blocksize,
smem
);
int maxBlocks = maxBlocksPerSM * numSMs;
maxBlocks -= maxBlocks % wkBlocksPerEnv; //round down to nearest multiple of wkBlocksPerEnv
printf("wkBlocksPerEnv %d, maxBlocks: %d\n", wkBlocksPerEnv, maxBlocks);
int* d_array;
unsigned char* d_synctemp;
cudaMalloc(&d_array, sizeof(int) * maxBlocks);
cudaMalloc(&d_synctemp, sizeof(unsigned char) * maxBlocks);
cudaMemset(d_synctemp, 0, sizeof(unsigned char) * maxBlocks);
kernel<<<maxBlocks, blocksize>>>(d_synctemp, d_array);
cudaFree(d_synctemp);
cudaFree(d_array);
return 0;
}
I like to do CUDA synchronization for multiple blocks.
你应该学会dis-like它。同步总是代价高昂的，即使实施得恰到好处，inter-core 同步更是如此。
if (threadIdx.x == 0){
// incrementing env_sync_block_count by 1
atomicAdd(&env_sync_block_count[kThisEnvId], 1);
while(env_sync_block_count[kThisEnvId] != wkBlocksPerEnv)
// OH NO!!
{
}
}
这很糟糕。使用此代码，每个块的第一个 warp 将执行 env_sync_block_count[kThisEnvId] 的重复读取。首先，正如@AbatorAbetor 提到的，您将面临 缓存不连贯 的问题，导致您的块在全局值长期更改后很可能从本地缓存中读取错误的值。
此外，您的块会占用多处理器。块将无限期地驻留并至少有一个活动扭曲。谁敢说将从他们的多处理器中逐出以安排额外的块来执行？如果我是 GPU，我不会让越来越多的活动块堆积起来。即使你没有死锁 - 你也会浪费很多时间。
现在，@AbatorAbetor 的回答通过限制网格大小避免了死锁。我想这行得通。但是除非你有充分的理由以这种方式编写你的内核 - 真正的解决方案是将你的算法分解成连续的内核（或者更好的是，弄清楚如何避免需要完全同步）。
一种mid-way方法是只有一些块通过同步点。你可以通过不等待来做到这一点，除非在某些情况下适用于非常有限数量的块（假设你有一个工作组 - 然后只有获得最后 K 个可能计数器值的块，等待）。
我喜欢为多个块做CUDA同步。并不是每个 __syncthreads() 可以轻松处理的块。
我看到有关于这个主题的现有讨论，例如 cuda block synchronization, and I like the simple solution brought up by @johan, ，本质上它使用 64 位计数器来跟踪同步块。
然而，我写了下面的代码试图完成类似的工作但是遇到了一个问题。这里我使用了术语 environment 以便此环境中的块 wkNumberEnvs 应该同步。它有一个柜台。我使用 atomicAdd() 来计算自己已经同步了多少个块，一旦同步块的数量 == wkBlocksPerEnv，我就知道所有块都已完成同步并且可以自由使用。但是，它有一个奇怪的结果，我不确定为什么。
问题出在这个while循环中。由于所有块的第一个线程都在执行 atomicAdd，因此有一个 while 循环来检查直到条件满足。但是我发现有些block会陷入死循环，不知道为什么最终不能满足条件？如果我在 *** I can print here 1 或 *** I can print here 2 中打印一些消息，就没有无限循环，一切都很完美。我没有看到明显的东西。
const int wkBlocksPerEnv = 2;
__device__ int env_sync_block_count[wkNumberEnvs];
__device__ void syncthreads_for_env(){
// sync threads for each block so all threads in this block finished the previous tasks
__syncthreads();
// sync threads for wkBlocksPerEnv blocks for each environment
if(wkBlocksPerEnv > 1){
const int kThisEnvId = get_env_scope_block_id(blockIdx.x);
if (threadIdx.x == 0){
// incrementing env_sync_block_count by 1
atomicAdd(&env_sync_block_count[kThisEnvId], 1);
// *** I can print here 1
while(env_sync_block_count[kThisEnvId] != wkBlocksPerEnv){
// *** I can print here 2
}
// Do the next job ...
}
}
原子值将进入全局内存，但在 while-loop 中您直接读取它并且它必须来自不会在线程之间自动同步的缓存（cache-coherence 仅由显式同步处理像 threadfence）。线程获得自己的同步，但其他线程可能看不到它。
即使您使用 threadfence，同一个 warp 中的线程也会 dead-lock 如果它们是第一个在任何其他线程更新它之前检查该值的线程，那么它们将永远等待。但应该与支持独立线程调度的最新 GPU 一起工作。
您的代码存在两个潜在问题。缓存和块调度。
缓存可以防止您在 while 循环期间观察更新的值。
如果您等待尚未安排的块更新，则块安排可能会导致 dead-lock。由于 CUDA 不保证调度块的特定顺序，因此防止此 dead-lock 的唯一方法是限制网格中的块数，以便所有块可以同时 运行。
以下代码显示了如何在避免上述问题的同时同步多个块。我从 CUDA-sample conjugateGradientMultiDeviceCG https://github.com/NVIDIA/cuda-samples/blob/master/Samples/4_CUDA_Libraries/conjugateGradientMultiDeviceCG/conjugateGradientMultiDeviceCG.cu#L186
在 pre-Volta 设备上，它使用易失性内存访问。 Volta 和后来使用 acquire/release 语义。 网格大小受查询设备属性限制。
#include <cassert>
#include <cstdio>
constexpr int wkBlocksPerEnv = 13;
__device__
int getEnv(int blockId){
return blockId / wkBlocksPerEnv;
}
__device__
int getRankInEnv(int blockId){
return blockId % wkBlocksPerEnv;
}
__device__
unsigned char load_arrived(unsigned char *arrived) {
#if __CUDA_ARCH__ < 700
return *(volatile unsigned char *)arrived;
#else
unsigned int result;
asm volatile("ld.acquire.gpu.global.u8 %0, [%1];"
: "=r"(result)
: "l"(arrived)
: "memory");
return result;
#endif
}
__device__
void store_arrived(unsigned char *arrived,
unsigned char val) {
#if __CUDA_ARCH__ < 700
*(volatile unsigned char *)arrived = val;
#else
unsigned int reg_val = val;
asm volatile(
"st.release.gpu.global.u8 [%1], %0;" ::"r"(reg_val) "l"(arrived)
: "memory");
// Avoids compiler warnings from unused variable val.
(void)(reg_val = reg_val);
#endif
}
#if 0
//wrong implementation which does not synchronize. to check that kernel assert does trigger without proper synchronization
__device__
void syncthreads_for_env(unsigned char* temp){
}
#else
//temp must have at least size sizeof(unsigned char) * total_number_of_blocks in grid
__device__
void syncthreads_for_env(unsigned char* temp){
__syncthreads();
const int env = getEnv(blockIdx.x);
const int blockInEnv = getRankInEnv(blockIdx.x);
unsigned char* const mytemp = temp + env * wkBlocksPerEnv;
if(threadIdx.x == 0){
if(blockInEnv == 0){
// Leader block waits for others to join and then releases them.
// Other blocks in env can arrive in any order, so the leader have to wait for
// all others.
for (int i = 0; i < wkBlocksPerEnv - 1; i++) {
while (load_arrived(&mytemp[i]) == 0)
;
}
for (int i = 0; i < wkBlocksPerEnv - 1; i++) {
store_arrived(&mytemp[i], 0);
}
__threadfence();
}else{
// Other blocks in env note their arrival and wait to be released.
store_arrived(&mytemp[blockInEnv - 1], 1);
while (load_arrived(&mytemp[blockInEnv - 1]) == 1)
;
}
}
__syncthreads();
}
#endif
__global__
void kernel(unsigned char* synctemp, int* array){
const int env = getEnv(blockIdx.x);
const int blockInEnv = getRankInEnv(blockIdx.x);
if(threadIdx.x == 0){
array[blockIdx.x] = 1;
}
syncthreads_for_env(synctemp);
if(threadIdx.x == 0){
int sum = 0;
for(int i = 0; i < wkBlocksPerEnv; i++){
sum += array[env * wkBlocksPerEnv + i];
}
assert(sum == wkBlocksPerEnv);
}
}
int main(){
const int smem = 0;
const int blocksize = 128;
int deviceId = 0;
int numSMs = 0;
int maxBlocksPerSM = 0;
cudaGetDevice(&deviceId);
cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, deviceId);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&maxBlocksPerSM,
kernel,
blocksize,
smem
);
int maxBlocks = maxBlocksPerSM * numSMs;
maxBlocks -= maxBlocks % wkBlocksPerEnv; //round down to nearest multiple of wkBlocksPerEnv
printf("wkBlocksPerEnv %d, maxBlocks: %d\n", wkBlocksPerEnv, maxBlocks);
int* d_array;
unsigned char* d_synctemp;
cudaMalloc(&d_array, sizeof(int) * maxBlocks);
cudaMalloc(&d_synctemp, sizeof(unsigned char) * maxBlocks);
cudaMemset(d_synctemp, 0, sizeof(unsigned char) * maxBlocks);
kernel<<<maxBlocks, blocksize>>>(d_synctemp, d_array);
cudaFree(d_synctemp);
cudaFree(d_array);
return 0;
}
I like to do CUDA synchronization for multiple blocks.
你应该学会dis-like它。同步总是代价高昂的，即使实施得恰到好处，inter-core 同步更是如此。
if (threadIdx.x == 0){
// incrementing env_sync_block_count by 1
atomicAdd(&env_sync_block_count[kThisEnvId], 1);
while(env_sync_block_count[kThisEnvId] != wkBlocksPerEnv)
// OH NO!!
{
}
}
这很糟糕。使用此代码，每个块的第一个 warp 将执行 env_sync_block_count[kThisEnvId] 的重复读取。首先，正如@AbatorAbetor 提到的，您将面临 缓存不连贯 的问题，导致您的块在全局值长期更改后很可能从本地缓存中读取错误的值。
此外，您的块会占用多处理器。块将无限期地驻留并至少有一个活动扭曲。谁敢说将从他们的多处理器中逐出以安排额外的块来执行？如果我是 GPU，我不会让越来越多的活动块堆积起来。即使你没有死锁 - 你也会浪费很多时间。
现在，@AbatorAbetor 的回答通过限制网格大小避免了死锁。我想这行得通。但是除非你有充分的理由以这种方式编写你的内核 - 真正的解决方案是将你的算法分解成连续的内核（或者更好的是，弄清楚如何避免需要完全同步）。
一种mid-way方法是只有一些块通过同步点。你可以通过不等待来做到这一点，除非在某些情况下适用于非常有限数量的块（假设你有一个工作组 - 然后只有获得最后 K 个可能计数器值的块，等待）。