矩陣未在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錯誤檢查添加到代碼中，則可以使用它（推薦）。