c - CUDA矩阵水平翻转中的无效参数

Question

当我们在 Linux 上的 CUDA Visual Profiler 下运行我们的 CUDA 应用程序时，我们遇到了问题。当我们创建一个新 session 并且工具包生成时间线时，如果矩阵大小非常大，则不会出现内核调用。我们的应用程序旨在水平翻转由随机数组成的矩阵，我们在命令行上传递它的大小。我们必须使用 cuda_visual_profiler 库的例程和 API 来占用 CPU 和 GPU 的时间，但我们不知道如何去做。

有人可以帮助我们吗？

谢谢大家

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

#include <thrust/system_error.h>
#include <thrust/system/cuda_error.h>
#include <sstream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/* START PROGRAM */

void inizializzaMatrice (int*,int,int);             // Initialize Matrix
void stampaMatrice (int*,int,int);                  // Print Matrix
void flipMatriceCPU (int*,int,int);                 // CPU Flipping
void confrontaMatrici (int*,int*,int,int);          // Equal Matrix
__global__ void flipMatriceGPU (int*,int*,int,int); // GPU Flipping

int main(int argn, char * argv[]){
    dim3 nBlocchi,nThreadPerBlocco;
    int M,N,flag;
    int *in_host, *out_host,*out_DeToHo;
    int *in_device, *out_device;
    long int size;
    printf("\n\n******************** RIFLESSIONE ORIZZONTALE DI UNA MATRICE ********************\n\n");
    if(argn<6 || atoi(argv[2])%2==0 ){
        if(argn<6)
            printf("Numero di parametri insufficiente!!!\n");

        else if(atoi(argv[2])%2==0)
            printf("Errore nell'utilizzo di %s. Il numero di colonne <N> deve essere dispari\n",argv[0]);

        printf("Uso corretto: %s <M> <N> <NumThreadPerBlocco.x> <NumThreadPerBlocco.y> <flag per la Stampa>\n", argv[0]);
        printf("Uso dei valori di default ... ...\n\n\n"); 
        nThreadPerBlocco.x=2; 
        nThreadPerBlocco.y=3;
        M=4; N=9; flag=1;
    }
    else {
        M=atoi(argv[1]); 
        N=atoi(argv[2]);
        nThreadPerBlocco.x=atoi(argv[3]);
        nThreadPerBlocco.y=atoi(argv[4]);
        flag=atoi(argv[5]); 
    }


    nBlocchi.x=M/nThreadPerBlocco.x+((M%nThreadPerBlocco.x)==0?0:1); 
    nBlocchi.y=N/nThreadPerBlocco.y+((N%nThreadPerBlocco.y)==0?0:1); 

    size=M*N*sizeof(int);

//stampa delle info sull'esecuzione del kernel
    printf("Numero di elementi = %d * %d\n",M, N);
    printf("Numero di thread per blocco = %d * %d\n", nThreadPerBlocco.x,nThreadPerBlocco.y); 
    printf("Numero di blocchi = %d * %d\n\n\n",nBlocchi.x,nBlocchi.y);

// Allocazione dati sull'host
    in_host=(int*)malloc(size);
    out_host=(int*)malloc(size);
    out_DeToHo=(int*)malloc(size);

// Allocazione dati dul device

    gpuErrchk( cudaMalloc((void**)&in_device,size) );
    gpuErrchk( cudaMalloc((void**)&out_device,size) );

// Inizializzazione dati sull'host

    inizializzaMatrice(in_host,M,N);

//cudaProfilerStart();

// Copia dei dati dall'host al device

    gpuErrchk(cudaMemcpy(in_device, in_host, size, cudaMemcpyHostToDevice));

// Invocazione del Kernel

    flipMatriceGPU<<<nBlocchi, nThreadPerBlocco, size>>>(in_device, out_device, N,M);

// Copia dei risultati dal device all'host
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );
    gpuErrchk( cudaMemcpy(out_DeToHo, out_device, size, cudaMemcpyDeviceToHost) );


// Flip Matrice CPU
    memcpy(out_host,in_host,size);
    flipMatriceCPU(out_host,M,N);

//cudaProfilerStop();

// Stampa Matrici

    if (flag==1){

        printf("Matrice di input:\n");
        //stampaMatrice(in_host, M, N);

        printf("Matrice di output host CPU:\n");
        //stampaMatrice(out_host, M, N);

        printf("Matrice di output device GPU:\n");
        stampaMatrice(out_DeToHo, M, N);

    }

    confrontaMatrici(out_host,out_DeToHo,M,N);
    printf("\n\n********************************************************************************\n\n");
    free(in_host);
    free(out_host);
    free(out_DeToHo);
    cudaFree(in_device);
    cudaFree(out_device);

    exit(0);
}






void inizializzaMatrice(int* matrice, int M, int N) {
    int i,j; for(i=0;i<M;i++)
    for(j=0;j<N;j++) matrice[i*N+j]=i*N+j;
}

void stampaMatrice(int*matrice, int M, int N) {
    int i,j; 
    for(i=0;i<M;i++) {
        for(j=0;j<N;j++)
            printf("%d\t", matrice[i*N+j]);
        printf("\n"); 
    }
}

void flipMatriceCPU(int *matrice, int row, int col){
    int i, j,tmp;
    for ( i = 0; i < row; i++ ) {
        for (  j = 0; j < col/2; j++ ) {
            tmp=matrice[col*i+j];
            matrice[col*i+j] = matrice[col*i+col-j-1];
            matrice[col*i+col-j-1] = tmp;

        }
    }
}

void confrontaMatrici(int* m1, int*m2, int M, int N) {
    int i, j; for(i=0;i<M;i++)
    for(j=0;j<N;j++) if(m1[i*N+j]!=m2[i*N+j]) {
        printf("I risultati dell'host e del device sono diversi.\n");
        return; 
    }
    if(i==M && j==N)
        printf("I risultati dell'host e del device coincidono.\n");
}

__global__ void flipMatriceGPU(int *in, int *out, int col, int row) {
    extern __shared__ int s_data[];
    int indexRow=threadIdx.x + blockIdx.x*blockDim.x; 
    int indexCol=threadIdx.y + blockIdx.y*blockDim.y; 
    int index=indexRow*col+indexCol;

    if(indexCol<col && indexRow<row){
        int index_data=blockDim.y-1-threadIdx.y+indexRow*col;

        s_data[index_data]=in[index];
        __syncthreads();

        int outOffset= blockDim.y*(gridDim.y-1-blockIdx.y);
        int outIndex= outOffset + threadIdx.y -(gridDim.y*blockDim.y - col) + indexRow*col;

        if(blockIdx.y==gridDim.y-1){
            outIndex+=gridDim.y*blockDim.y - col;
            out[outIndex]= s_data[(gridDim.y*blockDim.y - col)+(threadIdx.y+indexRow*col)];
        }

        else    
            out[outIndex]= s_data[threadIdx.y+indexRow*col];

    }
}

Answer 1

在 flipMatriceGPU __global__ 函数中，您正在动态分配大小为 size 的共享内存数组。当您考虑 100 x 131 大小的矩阵时，size 变得等于 52400，这超过了 的最大允许内存大小>48KB 每个流式多处理器。