库达异常答案

【问题标题】：Cuda Exceptions库达异常
【发布时间】：2011-11-30 23:44:05
【问题描述】：

我在 CUDA (FFT) 中做一些事情，但我不知道为什么它在调用内核函数时会产生异常。

所有包含和定义：

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <time.h>


#define CPU_ARRAY_SIZE 1024   // 1024, 2048, 4096 8192
#define GPU_ARRAY_SIZE 512   // 
#define THREAD_SIZE 16        // fixed
#define BLOCK_SIZE (GPU_ARRAY_SIZE/THREAD_SIZE)  // 32

#define PI 3.14

当我在 NVIDIA GTX480 中运行它时，我认为它可能是共享内存空间，尽管它似乎不是（因为有“一些”共享变量）。所以，我改变了 GPU_ARRAY_SIZE 来看看它是如何工作的，当我将它定义为 32、64、256、512 时，它给了我不同的结果（在 512 的情况下，它返回全零，我猜 CUDA 不能做任何事情 - 在其他情况下，它返回奇怪，因为我不知道它为什么会跳过 16 个单元格而不进行任何计算）。在大多数情况下，在我的 Microsoft Visual Studio 的“输出”窗口中，它会返回数十亿个异常，格式为“.exe 中 0x75b9b9bc 的第一次机会异常：Microsoft C++ 异常：内存位置的 cudaError_enum”。在您要求我调试之前，我无法调试它，因为 VS 不会对 VS 无法识别的文件执行此操作（例如 .cpp - 至少这个理论适用于我的情况）。各位大侠有什么想法吗： 1.为什么会产生异常？ 2. 为什么要计算，它应该对每个块中的每个单元格做什么，就在几个单元格内

我该如何解决这个问题...有什么想法吗？

内核函数：

__global__ void twiddle_factor(double *d_isub_matrix, double *d_osub_matrix)
{

    __shared__ double block[THREAD_SIZE][THREAD_SIZE];
    __shared__ double spectrum[THREAD_SIZE][THREAD_SIZE];
    __shared__ double sum_cos[THREAD_SIZE][THREAD_SIZE];  // declaring the shared sum_cos.. similarly for sum_sin
    __shared__ double sum_sin[THREAD_SIZE][THREAD_SIZE];
    __shared__ double local_cos[THREAD_SIZE][THREAD_SIZE];  // declaring the shared sum_cos.. similarly for sum_sin
    __shared__ double local_sin[THREAD_SIZE][THREAD_SIZE];

    unsigned int xIndex = threadIdx.x + blockIdx.x* blockDim.x;
    unsigned int yIndex = threadIdx.y + blockIdx.y* blockDim.y;


    int u;
    int x=0,y=0;

    int tx = threadIdx.x;
    int ty = threadIdx.y;

    double sum_sines=0.0,sum_cosines=0.0;

    double angle=(2*PI)/GPU_ARRAY_SIZE;       

    block[tx][ty] = d_isub_matrix[yIndex*GPU_ARRAY_SIZE+xIndex];

    __syncthreads();


    //for every column!

    for(u=0; u<THREAD_SIZE; u++)
    {

        /* All threads calculate its own sin and cos value. */
        local_sin[tx][ty] = block[tx][ty] * sin((angle*ty)*u);
        local_cos[tx][ty] = block[tx][ty] * cos((angle*ty)*u);


        /* Only one row is activate. The thread in row adds all element of its column. */
        if (ty == u) 
        {
            sum_sines   = 0.0;
            sum_cosines = 0.0;

            /* Access each column to add all elements of the column.*/
            for (y=0; y<THREAD_SIZE; y++)
            {
                sum_sines   += local_sin[tx][y];
                sum_cosines += local_cos[tx][y];
            }

            //if (sum_sines < 0) 
                //sum_sin[u][tx] = ((-1)*sum_sines)/GPU_ARRAY_SIZE;
            //else 
                sum_sin[u][tx] = sum_sines/GPU_ARRAY_SIZE;

            //if (sum_cosines < 0) 
                //sum_cos[u][tx] = ((-1)*sum_cosines)/GPU_ARRAY_SIZE;
            //else 
                sum_cos[u][tx] = sum_cosines/GPU_ARRAY_SIZE;

        }

        __syncthreads();
    }

    spectrum[tx][ty] = sqrt((double)pow(sum_sin[tx][ty],2) 
                               +(double)pow(sum_cos[tx][ty],2));
    __syncthreads();


    block[tx][ty] = spectrum[tx][ty];


    __syncthreads();


    //for every row!

    for(u=0; u<THREAD_SIZE; u++)
    {

        /* All threads calculate its own sin and cos value. */
        local_sin[tx][ty] = block[tx][ty] * sin((angle*ty)*u);
        local_cos[tx][ty] = block[tx][ty] * cos((angle*ty)*u);


        /* Only one column is activate. The thread in colum adds all element of its row. */
        if (tx == u) 
        {
            sum_sines   = 0.0;
            sum_cosines = 0.0;

            for (x=0; x<THREAD_SIZE; x++)
            {
                sum_sines   += local_sin[x][ty];
                sum_cosines += local_cos[x][ty];
            }

            //if (sum_sines < 0) 
                //sum_sin[ty][u] = ((-1)*sum_sines)/GPU_ARRAY_SIZE;
            //else 
                sum_sin[ty][u] = sum_sines/GPU_ARRAY_SIZE;

            //if (sum_cosines < 0) 
                //sum_cos[ty][u] = ((-1)*sum_cosines)/GPU_ARRAY_SIZE;
            //else 
                sum_cos[ty][u] = sum_cosines/GPU_ARRAY_SIZE;

        }

        __syncthreads();
    }

    spectrum[tx][ty] = sqrt((double)pow(sum_sin[tx][ty],2)+(double)pow(sum_cos[tx][ty],2));
    __syncthreads();


        /* Transpose! I think this is not necessary part. */

    d_osub_matrix[xIndex*GPU_ARRAY_SIZE + yIndex] =  spectrum[threadIdx.y][threadIdx.x];

    __syncthreads();
}

主要功能：

int main(int argc, char** argv)
{

    int i,j, w, h, sw, sh;

    int numSubblock = CPU_ARRAY_SIZE / GPU_ARRAY_SIZE;
        double *d_isub_matrix,*d_osub_matrix;

    double *big_matrix  = new double[CPU_ARRAY_SIZE*CPU_ARRAY_SIZE];
    double *big_matrix2 = new double[CPU_ARRAY_SIZE*CPU_ARRAY_SIZE];

    double *isub_matrix = new double[GPU_ARRAY_SIZE*GPU_ARRAY_SIZE];
    double *osub_matrix = new double[GPU_ARRAY_SIZE*GPU_ARRAY_SIZE];
    cudaEvent_t  start,stop;
    float elapsedtime;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);


    for (i=0; i<CPU_ARRAY_SIZE; i++)
    {
        for (j=0; j<CPU_ARRAY_SIZE; j++)
        big_matrix[i*CPU_ARRAY_SIZE + j] = rand();//i*CPU_ARRAY_SIZE + j;
    }   



    cudaEventRecord(start,0);


    //cudaMalloc((void**)&d_isub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float)*2);
    //cudaMalloc((void**)&d_osub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float)*2);

    for(i = 0; i < numSubblock; i++) 
    {
        for (j=0; j < numSubblock; j++) 
        {


        // start position of subarea of big array
        cudaMalloc((void**)&d_isub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float));
        cudaMalloc((void**)&d_osub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float));

        h = i*GPU_ARRAY_SIZE;

        w = j*GPU_ARRAY_SIZE;
        //printf("h = %d, w=%d",h,w);
        //system("PAUSE");

        // move subarea of big array into isub array.

        for (sh = 0; sh < GPU_ARRAY_SIZE; sh++)
        {
            for (sw = 0; sw <GPU_ARRAY_SIZE; sw++) 
            {
            isub_matrix[sh*GPU_ARRAY_SIZE+sw] = big_matrix[(h+sh)*CPU_ARRAY_SIZE + (w+sw)];

            }
        }



            cudaMemcpy(d_isub_matrix,isub_matrix,((GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float)),cudaMemcpyHostToDevice);

        //call the cuda kernel
        dim3 blocks(BLOCK_SIZE, BLOCK_SIZE);
        dim3 threads(THREAD_SIZE, THREAD_SIZE);

            twiddle_factor<<<blocks, threads>>>(d_isub_matrix,d_osub_matrix);

        cudaMemcpy(osub_matrix,d_osub_matrix,((GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float)),cudaMemcpyDeviceToHost);


        for (sh = 0; sh < GPU_ARRAY_SIZE; sh++)
        {
            for (sw = 0; sw <GPU_ARRAY_SIZE; sw++)
            {
                big_matrix2[(h+sh)*CPU_ARRAY_SIZE + (w+sw)] = osub_matrix[sh*GPU_ARRAY_SIZE+sw];
                printf(" sh %d  sw %d  %lf  \n", sh, sw, osub_matrix[sh*GPU_ARRAY_SIZE+sw]);

            }

        }
        printf("passei por aqui algumas vezes\n");
        cudaFree(d_osub_matrix);
        cudaFree(d_isub_matrix);

      }
    }
//  cudaFree(d_osub_matrix);
//  cudaFree(d_isub_matrix);

        //Stop the time
        cudaEventRecord(stop,0);
        cudaEventSynchronize(stop);
        cudaEventElapsedTime(&elapsedtime,start,stop);

    //showing the processing time
    printf("The processing time took... %fms to execute everything",elapsedtime);
    system("PAUSE");

        for (sh = 0; sh < CPU_ARRAY_SIZE; sh++)
        {
            for (sw = 0; sw <CPU_ARRAY_SIZE; sw++)
            {

                printf(" sh %d  sw %d  %lf  \n", sh, sw, big_matrix2[sh*CPU_ARRAY_SIZE+sw]);

            }
        }


    system("PAUSE");
    // I guess the result is "[1][0] = [1], [1][512] = [513], [513][0] = [524289], [513][512] = [524801]". 

}

【问题讨论】：

标签： cuda parallel-processing

【解决方案1】：

简单来说，问题可能而且应该是以下几行：

// start position of subarea of big array
  cudaMalloc((void**)&d_isub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float));
  cudaMalloc((void**)&d_osub_matrix,(GPU_ARRAY_SIZE*GPU_ARRAY_SIZE)*sizeof(float));

您只为 GPU 上的 double 值分配了很少的内存。您的子矩阵每个点分配 4 个字节，需要 8 个字节。

【讨论】：

好吧，看来你打对了！ :) 我乘以 2 你说我应该需要 8 个字节......那么至少它给了我答案！不错，克洛诺斯！ :) 感谢您的帮助！
嗯...这是一个很好的镜头...直到我开始将它们全部打印出来...从 sh=256 开始，它又开始给我零！ :(你对此有什么猜测吗？？（我又开始看到它，它重新开始给我一些好迹象，我猜在 sh=500 左右......非常奇怪！）
您是否将 cudaMemcpy 复制的日期数量更改为？他们的尺寸计算也使用 sizeof(float) 而不是 sizeof(double)
kronos，我刚刚做到了，直到现在（sha 已经通过了 520...）是好的！ :) 关于 sizeof 的事情还可以……它不会“干扰”这个过程。至少我是这样理解的，因为在内核中他们不接受双精度，只接受浮点数（这是我总是从 VS 的输出窗口中得到的）。 :P 无论如何...直到现在我可以说非常感谢！ :)
啊，是的，您使用的是 GTX 480，我一定已经监督了该信息；）很高兴看到它最终可以正常工作。