为什么这段 cuda 的并行前缀算法不会数据竞争

[rungo]

比如在 iter=1 这一轮的时候，同时会写入和读取 sum_buf[1]

naive_ker = SourceModule(""" 
                         __global__ void naive_prefix(double *vec, double *out) 
                         {  
                             __shared__ double sum_buf[1024];   
                             int tid = threadIdx.x;   
                             sum_buf[tid] = vec[tid]; 
                               
                             int iter = 1;  
                             for (int i=0; i < 10; i++)  
                             {  
                                 __syncthreads();  
                                 if (tid >= iter )  
                                 {  
                                     sum_buf[tid]= sum_buf[tid] + sum_buf[tid - iter];   
                                 }   
                                 iter *= 2;  
                             }  
                             __syncthreads(); 
                             out[tid] = sum_buf[tid];  
                             __syncthreads();  
                         }     
                         """) 
 naive_gpu = naive_ker.get_function("naive_prefix")

[r6cb]

你这都是写入到 sum_buf[tid]，每个线程 tid 不同，怎么会和其他线程冲突呢

[rungo]

@r6cb 比如 iter=1 时, tid=2 线程会读取 sum_buf[1] ，tid=1 线程会写入 sum_buf[1]

[flyqie]

好奇 AI 对这种问题的解释怎么样？

感觉现在 AI 应该对这种还算"聪明"？

[rungo]

@flyqie 问了几个 ai 不能统一，这段代码我从书上抄过来的