【发布时间】：2018-06-13 15:42:43
【问题描述】：
我有一个似乎有竞争条件的 CUDA 内核，我试图查明这种竞争条件的来源。我知道 cuda-memcheck 的“racecheck”工具，但是racecheck 告诉我使用小输入时没有危险，这实际上也与我自己的调查一致。对于大量输入，尽管racecheck 似乎需要永远（字面意思），所以我不能使用它。
简单解释一下，定义为__device__ 变量的一维向量d_mat_3d 用0 填充并加载到全局内存中。作为内核输入的两个大数组（d_A 和d_v）也在main 中定义并传递给内核。数组d_mat_3d，称为mat_2d 的一段被剪切，加载到共享内存中，并对其进行一些处理。然后，mat_2d 将被写回到全局内存上的d_mat_3d。
如此处所示，使用原子操作是因为不使用原子操作mat_2d 会遇到不同线程的竞争条件。
我想我仍然有某种竞争条件的原因是mat_3d 的结果每次都不同。
关于这种竞争条件可能来自哪里的任何想法？我可以采取任何步骤来清除它（除了工具竞赛检查）？如果您认为没有竞争条件的证据，您能解释一下为什么每次执行内核时都会为d_mat_3d 分配不同的值吗？
CUDA 9.0 / NVidia Titan Black / Ubuntu 16.04
#include <cstdlib>
#include <sstream>
#include <cstdio>
#include <cuda.h>
#include <cuda_runtime_api.h>
#define W 7 // fix limit for loops in kernel
#define SIZE 100 // defining matrix dimension
#define N_ELEM 10000 // no of elements in each vector
#define NTPB 1024 // no of threads per block
using namespace std;
__device__ float d_mat_3d[SIZE*SIZE*SIZE];
__global__ void cuda_kernel(float *d_A, float *d_v){
__shared__ float mat_2d[SIZE*SIZE]; // a 2D slice of 3D matrix d_mat_3d
unsigned int n = blockDim.x*blockIdx.x+threadIdx.x;
if(n >= N_ELEM)
return;
int x, y, z, i;
float r;
float A = d_A[n];
float v = d_v[n];
#pragma unroll
for(x=0; x<SIZE; x++){
// load mat_2d (on shared memory) using d_mat_3d (on global memory)
for(i=0; i<SIZE*SIZE; i++){
mat_2d[i] = d_mat_3d[i+x*SIZE*SIZE];
}
// sync threads as mat_2d is on shared memory
__syncthreads();
for(y=SIZE/2; y<SIZE/2+W; y++){
for(z=SIZE/2; z<SIZE/2+W; z++){
r = sqrt( pow(A,2) / v ); // no need to be in these loops. I know, but for my real case, it must be.
atomicAdd(&mat_2d[z+y*SIZE], r); // atomically add r
}
}
__syncthreads();
// write mat_2d (shared memory) back to mat_3d (global memory)
for(i=0; i<SIZE*SIZE; i++){
d_mat_3d[i+x*SIZE*SIZE] = mat_2d[i];
}
}
}
// this function writes h_mat_3d to disk.
void write_image(float *h_mat_3d){
ostringstream o_addToFile;
o_addToFile << "mat3d.bin";
FILE *pFile;
pFile = fopen(o_addToFile.str().c_str(), "wb");
for(int i=0; i<SIZE*SIZE*SIZE; i++){
fwrite(&h_mat_3d[i], sizeof(float), 1, pFile);
}
fclose (pFile);
}
int main(){
int i;
float *h_A = new float[N_ELEM]; // some large vector
float *h_v = new float[N_ELEM]; // some other large vector
float h_mat_3d[SIZE*SIZE*SIZE]; // will be filled w/ 0
float *d_A; // device variables
float *d_v;
for(i=0; i<N_ELEM; i++){
h_A[i] = 0.2f+(float)i/N_ELEM; // fill out with some calculations
h_v[i] = 0.5f+2.f*i/N_ELEM;
}
for(i=0; i<SIZE*SIZE*SIZE; i++){
h_mat_3d[i] = 0.f; // fill h_mat_3d with 0
}
cudaMalloc((void **)&d_A, sizeof(float)*N_ELEM); // allocate variables on device
cudaMalloc((void **)&d_v, sizeof(float)*N_ELEM);
cudaMemcpy(d_A, h_A, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice); // copy from host to device
cudaMemcpy(d_v, h_v, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(d_mat_3d, &h_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // copy h_mat_3d to device
cuda_kernel<<<(N_ELEM+NTPB-1)/NTPB,NTPB>>>(d_A, d_v); // execute kernel
cudaMemcpyFromSymbol(h_mat_3d, d_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // write it back to h_mat_3d
write_image(h_mat_3d); // write h_mat_3d to disk for checking
cudaFree(d_A); // free memory
cudaFree(d_v);
delete [] h_A;
delete [] h_v;
return 0;
}
【问题讨论】：
-
实际上对于此类问题，SO 希望您提供minimal reproducible example。阅读第 1 项here。注意“必须”这个词的用法。如果您的声明是正确的，则不需要很多额外的代码行来构建一个完整的程序来证明您的声明。
标签： c++ cuda race-condition