簡單的CUDA測試總是失敗，出現「非法內存訪問」錯誤

Question

如果我運行這個程序，我得到「在第48行matrixMulti.cu出現非法內存訪問」錯誤。我搜索並嘗試了很多。所以我希望有人能幫助我。 （數組，devarray，N * N * sizeof（int），cudaMemcpyDeviceToHost））;

該程序只是爲了進入CUDA。我試圖實現矩陣乘法。

#include <iostream> 
#include<cuda.h> 
#include <stdio.h> 

using namespace std; 

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__)) 
void printVec(int** a, int n); 

static void HandleError(cudaError_t err, const char *file, int line) 
{ 
    if (err != cudaSuccess) 
    { 
    printf("%s in %s at line %d\n", cudaGetErrorString(err), 
      file, line); 
    exit(EXIT_FAILURE); 
    } 
} 

void checkCUDAError(const char *msg) 
{ 
    cudaError_t err = cudaGetLastError(); 
    if(cudaSuccess != err) 
    { 
     fprintf(stderr, "Cuda error: %s: %s.\n", msg, 
           cudaGetErrorString(err)); 
     exit(EXIT_FAILURE); 
    }       
} 
__global__ void MatrixMulti(int** a, int** b) { 
    b[0][0]=4; 
} 

int main() { 
    int N =10; 
    int** array, **devarray; 
    array = new int*[N]; 

    for(int i = 0; i < N; i++) { 
     array[i] = new int[N]; 
    } 

    HANDLE_ERROR (cudaMalloc((void**)&devarray, N*N*sizeof(int))); 
    HANDLE_ERROR (cudaMemcpy(devarray, array, N*N*sizeof(int), cudaMemcpyHostToDevice)); 
    MatrixMulti<<<1,1>>>(array,devarray); 
    HANDLE_ERROR (cudaMemcpy(array, devarray, N*N*sizeof(int), cudaMemcpyDeviceToHost)); 
    HANDLE_ERROR (cudaFree(devarray)); 
    printVec(array,N); 

    return 0; 
} 

void printVec(int** a , int n) { 
    for(int i =0 ; i < n; i++) { 
     for (int j = 0; j <n; j++) { 
     cout<< a[i][j] <<" "; 
     }  
     cout<<" "<<endl;  
    } 
} 

Answer 1

一般來說，您分配和複製一個雙下標C數組的方法將不起作用。 cudaMemcpy預計平面，連續分配，單指針，單下標數組。

由於這種混亂的結果，指針被傳遞到你的內核（int** a, int** b）無法正常（安全）解除引用兩次：

b[0][0]=4; 

當您嘗試做上述內核代碼，你會得到一個非法的內存訪問，因爲你沒有正確地分配設備上的指針指針風格分配。

如果你用cuda-memcheck運行你的代碼，你會在內核代碼中得到另一個非法內存訪問的跡象。

在這些情況下，通常的建議是將2D數組「扁平化」爲單維，並使用適當的指針或索引算術來模擬2D訪問。它可能可能分配二維數組（即雙下標，雙指針），但它是相當涉及（部分由於需要「深層複製」）。如果您想了解更多，只需在CUDA 2D array的右上角搜索即可。

這裏有一個版本的代碼具有陣列平整的設備端陣列：

$ cat t60.cu 
#include <iostream> 
#include <cuda.h> 
#include <stdio.h> 

using namespace std; 

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__)) 
void printVec(int** a, int n); 

static void HandleError(cudaError_t err, const char *file, int line) 
{ 
    if (err != cudaSuccess) 
    { 
    printf("%s in %s at line %d\n", cudaGetErrorString(err), 
      file, line); 
    exit(EXIT_FAILURE); 
    } 
} 

void checkCUDAError(const char *msg) 
{ 
    cudaError_t err = cudaGetLastError(); 
    if(cudaSuccess != err) 
    { 
     fprintf(stderr, "Cuda error: %s: %s.\n", msg, 
           cudaGetErrorString(err)); 
     exit(EXIT_FAILURE); 
    } 
} 

__global__ void MatrixMulti(int* b, unsigned n) { 
    for (int row = 0; row < n; row++) 
     for (int col=0; col < n; col++) 
    b[(row*n)+col]=col; //simulate 2D access in kernel code 
} 

int main() { 
    int N =10; 
    int** array, *devarray; // flatten device-side array 
    array = new int*[N]; 
    array[0] = new int[N*N]; // host allocation needs to be contiguous 
    for (int i = 1; i < N; i++) array[i] = array[i-1]+N; //2D on top of contiguous allocation 

    HANDLE_ERROR (cudaMalloc((void**)&devarray, N*N*sizeof(int))); 
    HANDLE_ERROR (cudaMemcpy(devarray, array[0], N*N*sizeof(int), cudaMemcpyHostToDevice)); 
    MatrixMulti<<<1,1>>>(devarray, N); 
    HANDLE_ERROR (cudaMemcpy(array[0], devarray, N*N*sizeof(int), cudaMemcpyDeviceToHost)); 
    HANDLE_ERROR (cudaFree(devarray)); 
    printVec(array,N); 

    return 0; 
} 

void printVec(int** a , int n) { 
    for(int i =0 ; i < n; i++) { 
     for (int j = 0; j <n; j++) { 
     cout<< a[i][j] <<" "; 
     } 
     cout<<" "<<endl; 
    } 
} 
$ nvcc -arch=sm_20 -o t60 t60.cu 
$ ./t60 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
0 1 2 3 4 5 6 7 8 9 
$