简单的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 
$