Я разрабатываю матричное умножение CUDA, но я сделал некоторые модификации, чтобы посмотреть, как они влияют на производительность.
Я пытаюсь наблюдать поведение (и я измеряю изменения во времени событий графического процессора) простого ядра умножения матриц. Но я тестирую его в двух разных условиях:
У меня есть количество матриц (скажем, matN) для A, B и C, затем я передаю (H2D) одну матрицу для A, одну для B за раз, а затем умножаю их, чтобы перенести обратно (D2H) один С;
У меня есть matN для A, B и C, но я перевожу> 1 (скажем, chunk) матрицы одновременно для A и B, точно выполняю chunk умножения и возвращаю chunk результат матрицы C.
В первом случае (chunk = 1) все работает как положено, но во втором случае (chunk > 1) я получаю, что некоторые из C верны, а другие неправильны.
Но если я поставлю cudaDeviceSynchronize() после cudaMemcpyAsync, все результаты, которые я получу, верны.
Вот часть кода, выполняющая то, что я только что описал выше:
/**** main.cpp ****/
int chunk = matN/iters;
#ifdef LOWPAR
GRIDx= 1;
GRIDy= 1;
label="LOW";
#else
int sizeX = M;
int sizeY = N;
GRIDx = ceil((sizeX)/BLOCK);
GRIDy = ceil((sizeY)/BLOCK);
label="";
#endif
const int bytesA = M*K*sizeof(float);
const int bytesB = K*N*sizeof(float);
const int bytesC = M*N*sizeof(float);
//device mem allocation
float *Ad, *Bd, *Cd;
gpuErrchk( cudaMalloc((void **)&Ad, bytesA*chunk) );
gpuErrchk( cudaMalloc((void **)&Bd, bytesB*chunk) );
gpuErrchk( cudaMalloc((void **)&Cd, bytesC*chunk) );
//host pinned mem allocation
float *A, *B, *C;
gpuErrchk( cudaMallocHost((void **)&A, bytesA*matN) );
gpuErrchk( cudaMallocHost((void **)&B, bytesB*matN) );
gpuErrchk( cudaMallocHost((void **)&C, bytesC*matN) );
//host data init
for(int i=0; i<matN; ++i){
randomMatrix(M, K, A+(i*M*K));
randomMatrix(K, N, B+(i*K*N));
}
//event start
createAndStartEvent(&startEvent, &stopEvent);
if (square)
{
label += "SQUARE";
int size = N*N;
for (int i = 0; i < iters; ++i) {
int j = i%nStream;
int idx = i*size*chunk;
newSquareMatMulKer(A+idx, B+idx, C+idx, Ad, Bd, Cd, N, chunk, stream[j]);
}
}
else {
...
}
msTot = endEvent(&startEvent, &stopEvent);
#ifdef MEASURES
printMeasures(square, label, msTot, millis.count(), matN, iters, devId);
#else
float *_A, *_B, *_C, *tmpC;
tmpC = (float *)calloc(1,bytesC*chunk);
for (int s=0; s<matN; ++s)
{
_A = A+(s*M*K);
_B = B+(s*K*N);
_C = C+(s*M*N);
memset(tmpC, 0, bytesC*chunk);
hostMatMul(_A, _B, tmpC, M, K, N);
checkMatEquality(_C, tmpC, M, N);
}
#endif
/**** matmul.cu ****/
__global__ void squareMatMulKernel(float* A, float* B, float* C, int N, int chunk) {
int ROW = blockIdx.x*blockDim.x+threadIdx.x;
int COL = blockIdx.y*blockDim.y+threadIdx.y;
if (ROW<N && COL<N) {
int size=N*N;
int offs = 0;
float tmpSum=0.0f;
for (int s=0; s<chunk; ++s)
{
offs = s*size;
tmpSum = 0.0f;
for (int i = 0; i < N; ++i) {
tmpSum += A[offs+(ROW*N)+i] * B[offs+(i*N)+COL];
}
C[offs+(ROW*N)+COL] = tmpSum;
}
}
return ;
}
void newSquareMatMulKer(float *A, float *B, float *C, float *Ad, float *Bd, float *Cd,
int n, int chunk, cudaStream_t strm)
{
int size = n*n;
int bytesMat = size*sizeof(float);
dim3 dimBlock(BLOCK,BLOCK,1);
dim3 dimGrid(GRIDx, GRIDy,1);
gpuErrchk( cudaMemcpyAsync(Ad, A, bytesMat*chunk, cudaMemcpyHostToDevice, strm) );
gpuErrchk( cudaMemcpyAsync(Bd, B, bytesMat*chunk, cudaMemcpyHostToDevice, strm) );
#ifdef LOWPAR
squareMatMulGridStrideKer<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);
#else
squareMatMulKernel<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);
#endif
squareMatMulKernel<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);
gpuErrchk( cudaMemcpyAsync( C, Cd, bytesMat*chunk, cudaMemcpyDeviceToHost, strm) );
cudaDeviceSynchronize();
^ ^ ^ ^ ^ ^
}
Я попытался отладить с помощью cuda-gdb, но ничего странного не обнаружилось, gpuErrchk не выдает никаких ошибок в вызовах API CUDA.
Я также запускаю код с помощью memcheck, как в случае с cudaDeviceSynchronize, так и без него, и в обоих случаях я не получаю ошибки.
Я думаю, что могу утверждать, что это проблема синхронизации, но я не могу понять, в чем причина этого.
Может кто-то определить, где я иду не так?
Также очень ценятся и другие советы по стилю кода.