矩阵未在CUDA中成功从设备复制回主机

Question

我是cuda新手。我编写了一个内核来创建维度大小为Xsize的标识矩阵（GPUsetIdentity）。在函数GPUfunctioncall的内部，我调用了我的内核。单位矩阵应该存储在dDataInv中。但是，当我将它复制回dataOut sizexsize时，所有值都为零。我知道，我在某个地方做了一些非常愚蠢的事情，但无法得到它，如果任何人都可以指出我的错误，我对cuda很陌生。谢谢。

#include <stdio.h> 
#include <malloc.h> 
#include <memory.h> 
#include <math.h> 
#include <stdlib.h> 
#include <iostream> 
#include <stdlib.h> 
#include <string> 
#include <fstream> 
#include <iterator> 
#include <sstream> 
#include <vector> 
#include <cstring> 
#include <cstdlib> 
#include <ctime> 
#include <stdlib.h> 
#include <cuda_runtime.h> 
#include "cuda.h" 

#define BLOCKSIZE 16 


using namespace std; 

__global__ void GPUsetIdentity (float* matrix, int width) 

{ 
     int tx = threadIdx.x; 
     int bx = blockIdx.x; 
     int offset = bx * BLOCKSIZE + tx; 
     matrix[offset + width * offset] = 1; 

} 


void print_matrix_host(float* A , int nr_rows_A, int nr_cols_A) { 

     for(int i = 0; i < nr_rows_A; ++i){ 
       for(int j = 0; j < nr_cols_A; ++j){ 
         std::cout << A[i * nr_rows_A + j ] << " "; 
       } 
       std::cout << std::endl; 

     } 
     std::cout << std::endl; 
} 

int GPUfunctioncall (float* hDataOut, int size){ 

     float *dDataInv; 


     cudaMalloc ((void **) &dDataInv, size); 
     cudaMemset ((void *) dDataInv, 0, size); 



     dim3 idyThreads (BLOCKSIZE); 
     dim3 idyBlocks (size/BLOCKSIZE); 


     GPUsetIdentity <<< idyBlocks, idyThreads >>> (dDataInv, size); 
     cudaThreadSynchronize(); 

     cudaMemcpy ((void *) hDataOut, (void *) dDataInv, size, cudaMemcpyDeviceToHost); 
     cudaFree (dDataInv); 

     return 0; 

} 

int main() 

{ 
     int size = 4; 
     float* dataOut; 

     dataOut = new float[size*size]; 

     GPUfunctioncall(dataOut, size); 
     print_matrix_host(dataOut, size, size); 


} 

Answer 1

您有一个CUDA代码的麻烦任何时候，它的使用proper cuda error checking好的做法。您也可以使用cuda-memcheck运行您的代码以快速阅读是否有任何错误。

使用这些方法中的任何一种，都会在内核启动时发现“无效配置错误”。这通常意味着<<< >>>语法中的参数不正确。当遇到这种类型的错误时，只需打印出这些值就可以指出问题所在。

在你的情况，这行代码：

dim3 idyBlocks (size/BLOCKSIZE); 

结果为idyBlocks一个的0值时size是4和BLOCKSIZE为16那么，你是请求内核启动的0块这是违法的。因此，你的内核不运行，你的结果不是你所期望的。

有很多种方法可以解决这个问题，其中许多方法涉及到检测到这种情况并在size不能被BLOCKSIZE整除时添加“额外的块”。使用这种方法，我们可能会启动“额外的线程”，所以我们必须在内核中包含一个“线程检查”，以防止这些额外的线程做任何事情（比如访问数组越界）。为此，我们经常需要知道内核中的预期大小，我们可以将此值作为额外的内核参数传递。

您在处理设备变量时也发生了一些错误。以下代码：

dataOut = new float[size*size]; 

为尺寸为size的方阵分配足够的空间。但是，下面的代码：

cudaMalloc ((void **) &dDataInv, size); 

只分配了size字节足够的空间。您希望size*size*sizeof(float)而不是size此处，并且您希望它在以下cudaMemset和cudaMemcpy操作。 cudaMalloc,cudaMemset和cudaMemcpy需要字节中的大小参数，就像malloc,memset和memcpy一样。在使用cudaMemset和cudaMemcpy时也会发现此错误。

下面的代码有这些修改，并似乎为我正常工作：

$ cat t580.cu 
#include <stdio.h> 
#include <stdlib.h> 
#include <iostream> 

#define BLOCKSIZE 16 


using namespace std; 

__global__ void GPUsetIdentity (float* matrix, int width, int size) 

{ 
     int tx = threadIdx.x; 
     int bx = blockIdx.x; 
     int offset = bx * BLOCKSIZE + tx; 
     if (tx < size) 
      matrix[offset + width * offset] = 1; 

} 


void print_matrix_host(float* A , int nr_rows_A, int nr_cols_A) { 

     for(int i = 0; i < nr_rows_A; ++i){ 
       for(int j = 0; j < nr_cols_A; ++j){ 
         std::cout << A[i * nr_rows_A + j ] << " "; 
       } 
       std::cout << std::endl; 

     } 
     std::cout << std::endl; 
} 

int GPUfunctioncall (float* hDataOut, int size){ 

     float *dDataInv; 


     cudaMalloc ((void **) &dDataInv, size*size*sizeof(float)); 
     cudaMemset ((void *) dDataInv, 0, size*size*sizeof(float)); 



     dim3 idyThreads (BLOCKSIZE); 
     int num_blocks = size/BLOCKSIZE + (size%BLOCKSIZE)?1:0; 
     dim3 idyBlocks (num_blocks); 


     GPUsetIdentity <<< idyBlocks, idyThreads >>> (dDataInv, size, size); 
     cudaThreadSynchronize(); 

     cudaMemcpy ((void *) hDataOut, (void *) dDataInv, size*size*sizeof(float), cudaMemcpyDeviceToHost); 
     cudaFree (dDataInv); 

     return 0; 

} 

int main() 

{ 
     int size = 4; 
     float* dataOut; 

     dataOut = new float[size*size]; 

     GPUfunctioncall(dataOut, size); 
     print_matrix_host(dataOut, size, size); 


} 
$ nvcc -arch=sm_20 -o t580 t580.cu 
$ cuda-memcheck ./t580 
========= CUDA-MEMCHECK 
1 0 0 0 
0 1 0 0 
0 0 1 0 
0 0 0 1 

========= ERROR SUMMARY: 0 errors 
$ 

注意，这可能是多余的传递size两次内核。对于这个特定的例子，我们可以很容易地使用width参数来做我们的内核“线程检查”。但出于教育目的，我选择将其作为单独的参数来调用它，因为在一般情况下，您通常会将它作为单独的参数传递给您编写的其他内核。

最后，请注意，cudaThreadSynchronize()是已弃用，应替换为cudaDeviceSynchronize()。在这个特殊的例子中，实际上它们是必需的，因为下一个cudaMemcpy操作将强制执行同一种类型的同步，但是如果您决定将cuda错误检查添加到代码中，则可以使用它（推荐）。