CUDA_SAFE_CALL: Es wurde ein unzulässiger Speicherzugriff festgestellt

Question

Ich versuche eine einfache Matrixmultiplikation auf CUDA durchzuführen. Ich weiß Arrays können abgeflacht werden, um es an das Gerät zu übergeben. Allerdings verwende ich cudaMallocPitch und cudaMemcpy2d, um die Multiplikation durchzuführen. Beim Ausführen des folgenden Codes bekomme ich einen Fehler "illegaler Speicher wurde gefunden", wenn ich versuche, das Ergebnis auf den Host zu kopieren. Ich schätze jeden Hinweis, wo ich falsch liege. Vielen Dank!

weights-erste Matrix, dim: 30x784

Input- zweite Matrix, dim: 784x100

results_d - Ergebnis auf dem Gerät (GPU)

Ergebnis - Ergebnis auf dem Host kopiert

#include <stdio.h> 
#include <math.h> 
#include <cstdio> 
#include <cstdlib> 

#define CUDA_SAFE_CALL(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"CUDA_SAFE_CALL: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

__global__ void MatrixMulKernel(double *input,double *weights,double *results_d,size_t in_pitch,size_t w1_pitch,size_t result_pitch) 
{ 
int row = threadIdx.x; 
int col= threadIdx.y; 
double value; 
double *result_matrix; 

result_matrix = ((double*)((char*)results_d + row*result_pitch + col)); 


printf("%d",threadIdx); 

for(int i =0 ; i < in_pitch ; i++) 

{ 

double *element1 = ((double*)((char*)input + row*in_pitch) + i) ; 
double *element2 = ((double*)((char*)weights + i*w1_pitch) + col); 

value =+ (*element1) * (*element2); 

} 

*result_matrix = value; 

} 





int main() 
{ 

static double arr1[30][784]; 
static double arr2[784][100]; 
static double result[30][100]; 



for (int i = 0 ; i < 30; i++) 

{ 
for(int j =0;j <784 ; j ++) 
arr1[i][j] = 5; 

} 

for (int i =0 ; i < 784; i ++) 
{ 

for(int j=0;j < 100 ; j++) 
arr2[i][j] = 3; 

} 



double *input; 
double *weights; 
double *results_d; 

size_t in_pitch,w1_pitch,result_pitch; 



//allocating memory in GPU for 2 inputs and result 
CUDA_SAFE_CALL(cudaMallocPitch((void**)&input,&in_pitch,100*sizeof(double),784)); 
CUDA_SAFE_CALL(cudaMallocPitch((void**)&weights,&w1_pitch,784*sizeof(double),30)); 
CUDA_SAFE_CALL(cudaMallocPitch((void**)&results_d,&result_pitch,100*sizeof(double),30)); 

//Copy matrix from host to device 
CUDA_SAFE_CALL(cudaMemcpy2D(input,in_pitch,arr2,100*sizeof(double),100*sizeof(double),784,cudaMemcpyHostToDevice)); 
CUDA_SAFE_CALL(cudaMemcpy2D(weights,w1_pitch,arr1,784*sizeof(double),784*sizeof(double),30,cudaMemcpyHostToDevice)); 
CUDA_SAFE_CALL(cudaMemcpy2D(results_d,result_pitch,result,100*sizeof(double),100*sizeof(double),30,cudaMemcpyHostToDevice)); 


//using GPU 


    dim3 dimGrid(1,1,1); 
    dim3 dimBlock(32,32,1); 
    printf("before kernel fucntion"); 
    MatrixMulKernel<<<dimGrid, dimBlock>>>(input, weights,results_d,in_pitch,w1_pitch,result_pitch);  
    printf("after kernel fucntion"); 
    cudaThreadSynchronize(); 

//copying back to host 
CUDA_SAFE_CALL(cudaMemcpy2D(result,result_pitch,results_d,100*sizeof(double),100*sizeof(double),30,cudaMemcpyDeviceToHost)); 


//printing and seeing whether the result matrix has been updated  
for (int i =0 ; i < 100; i ++) 
{ 

for(int j=0;j < 30 ; j++) 
{ 
printf("%f",result); 

} 
printf("\n"); 
} 

CUDA_SAFE_CALL(cudaFree(input)); 
CUDA_SAFE_CALL(cudaFree(weights)); 
CUDA_SAFE_CALL(cudaFree(results_d)); 


return 0; 
} 

Answer 1

Es gibt eine Reihe von Fehlern in diesem Code. Vor allem ist es nicht klar, dass die Zuteilung von Pitch hier einen Vorteil bringt. Zweitens sollten Sie CUBLAS verwenden, wenn Sie eine schnelle Matrix-Multiplikation wünschen.

Probleme:

1. Sie scheinen nicht aufgeschlagen Zuweisungen zu verstehen. Der zurückgegebene Wert pitch ist ein Wert in Bytes. Sie können das nicht sinnvoll für einen Schleifenindex für die Matrix-Multiplikation verwenden. Der Wert pitch entspricht der Gesamtbreite der Tonhöhenzuweisung. Es entspricht nicht dem gültigen Datenbereich. Dazu sollten Sie die entsprechende Matrixdimension verwenden.

2. Ihr Code führt keine Matrixmultiplikation über den gesamten Matrixbereich durch. Sie erstellen nur einen einzelnen Block von 32x32-Threads, aber Sie benötigen genügend Blöcke/Threads, um den gesamten Matrixbereich abzudecken. Dies erfordert Änderungen an den Grid-Dimensionen, Weiterleitung von Matrix-Dimensionen an Ihren Kernel sowie eine "Thread-Überprüfung" in Ihrem Kernel, um den Zugriff außerhalb der Grenzen zu verhindern.

3. Dieses Konstrukt für das Geneigte Zugang ist nicht korrekt:

   result_matrix = ((double*)((char*)results_d + row*result_pitch + col)); 
   

   es nicht die anderen Konstruktionen Sie für die 2 Eingangsmatrizen haben übereinstimmt, hat es eine fehl am Platze schließende Klammer.

4. Sie haben das Gefühl Ihrer zwei Eingangsmatrizen umgekehrt. Sie indexieren in die input Matrix, als wäre es die weight Matrix und umgekehrt. Wir müssen den Sinn von row, column und i austauschen, damit diese den tatsächlichen Matrixdimensionen entsprechen.

5. Ihre letzte cudaMemcpy2D Betrieb hat die Tonhöhenwerte umgekehrt:

cudaMemcpy2D(result,result_pitch,results_d,100*sizeof(double),100*sizeof(double),30,cudaMemcpyDeviceToHost)

    ^^^^^     ^^^^^ 

6. Sie haben vergessen, auf Null Ihre Schleife Summenvariable zu initialisieren :

   double value; 
   

7. Ich weiß nicht, was Sie hier soll, sollte es += nicht =+ sein:

Der folgende Code hat diese Probleme behoben, und scheint für mich ohne Fehler zu laufen

value =+ ... 

:

$ cat t104.cu 
#include <stdio.h> 
#include <math.h> 
#include <cstdio> 
#include <cstdlib> 

const int d1 = 30; 
const int d2 = 784; 
const int d3 = 100; 

double arr1[d1][d2]; 
double arr2[d2][d3]; 
double result[d1][d3]; 


#define CUDA_SAFE_CALL(ans) { gpuAssert((ans), __FILE__, __LINE__); } 

inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"CUDA_SAFE_CALL: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

__global__ void MatrixMulKernel(double *input,double *weights,double *results_d,size_t in_pitch,size_t w1_pitch,size_t result_pitch, int dim, int rrow, int rcol) 
{ 
    int col = threadIdx.x + blockDim.x*blockIdx.x; 
    int row= threadIdx.y + blockDim.y*blockIdx.y; 

    if ((row >= rrow) || (col >= rcol)) return; 

    double value = 0; 
    double *result_matrix; 

    result_matrix = ((double*)((char*)results_d + row*result_pitch) + col); 

    for(int i =0 ; i < dim ; i++) 

    { 

    double *element1 = ((double*)((char*)input + i*in_pitch) + col) ; 
    double *element2 = ((double*)((char*)weights + row*w1_pitch) + i); 

    value += (*element1) * (*element2); 

    } 

    *result_matrix = value; 

} 





int main() 
{ 


    for (int i = 0 ; i < d1; i++) 

    { 
    for(int j =0;j <d2 ; j ++) 
     arr1[i][j] = 5; 

    } 

    for (int i =0 ; i < d2; i ++) 
    { 

    for(int j=0;j < d3 ; j++) 
     arr2[i][j] = 3; 

    } 



    double *input; 
    double *weights; 
    double *results_d; 

    size_t in_pitch,w1_pitch,result_pitch; 



//allocating memory in GPU for 2 inputs and result 
    CUDA_SAFE_CALL(cudaMallocPitch((void**)&input,&in_pitch,d3*sizeof(double),d2)); 
    CUDA_SAFE_CALL(cudaMallocPitch((void**)&weights,&w1_pitch,d2*sizeof(double),d1)); 
    CUDA_SAFE_CALL(cudaMallocPitch((void**)&results_d,&result_pitch,d3*sizeof(double),d1)); 

//Copy matrix from host to device 
    CUDA_SAFE_CALL(cudaMemcpy2D(input,in_pitch,arr2,d3*sizeof(double),d3*sizeof(double),d2,cudaMemcpyHostToDevice)); 
    CUDA_SAFE_CALL(cudaMemcpy2D(weights,w1_pitch,arr1,d2*sizeof(double),d2*sizeof(double),d1,cudaMemcpyHostToDevice)); 
    CUDA_SAFE_CALL(cudaMemcpy2D(results_d,result_pitch,result,d3*sizeof(double),d3*sizeof(double),d1,cudaMemcpyHostToDevice)); 


//using GPU 


    dim3 dimBlock(32,32,1); 
    dim3 dimGrid(((d3+dimBlock.x-1)/dimBlock.x),((d1+dimBlock.y-1)/dimBlock.y),1); 
    MatrixMulKernel<<<dimGrid, dimBlock>>>(input, weights,results_d,in_pitch,w1_pitch,result_pitch, d2, d1, d3); 

//copying back to host 
    CUDA_SAFE_CALL(cudaMemcpy2D(result,d3*sizeof(double),results_d,result_pitch,d3*sizeof(double),d1,cudaMemcpyDeviceToHost)); 


//printing and seeing whether the result matrix has been updated 
    for (int i =0 ; i < d3; i ++) 
    { 

    for(int j=0;j < d1 ; j++) 
    { 
     printf("%f", result[j][i]); 

    } 
    printf("\n"); 
    } 

    CUDA_SAFE_CALL(cudaFree(input)); 
    CUDA_SAFE_CALL(cudaFree(weights)); 
    CUDA_SAFE_CALL(cudaFree(results_d)); 


    return 0; 
} 
$ nvcc -arch=sm_61 -o t104 t104.cu 
$