1

我刚开始学习 CUDA，只知道非常基本的东西。我正在尝试开发一个使用矩阵乘法方法执行 8x8 DCT 的 CUDA 程序。计算 8x8 DCT 系数矩阵 D，然后 DCT 变换为 D'AD。每个线程计算 1 个数据点，每个块为 8x8。我编写了一个顺序 DCT 并比较了输出文件中的结果。

这就是问题所在。当块数为 1xN 时，一切正常。当块数为 MxN 时（M 是任何大于 1 的数），内核给出错误的结果。我认为问题应该是我的块索引，但我找不到问题。

有人可以提供一些帮助吗？我知道这是一个非常基本的程序，但我真的需要它。任何意见将不胜感激！提前致谢！

#include <stdio.h>
#include <stdlib.h>
#include "types.h"
#include "cuda.h"

static int DCT_bases[64]= {2896, 2896, 2896, 2896, 2896, 2896, 2896, 2896,
    4017, 3406, 2276, 799, -799, -2276, -3406, -4017,
    3784, 1568, -1568, -3784, -3784, -1568, 1568, 3784,
    3406, -799, -4017, -2276, 2276, 4017, 799, -3406, 
    2896, -2896, -2896, 2896, 2896, -2896, -2896, 2896,
    2276, -4017, 799, 3406, -3406, -799, 4017, -2276, 
    1568, -3784, 3784, -1568, -1568, 3784, -3784, 1568,
    799, -2276, 3406, -4017, 4017, -3406, 2276, -799 };

__device__ __constant__ int dDCT_bases[64];

__global__ void cudaDCT2D(int *src, int width) {

    int i = blockIdx.y * blockDim.y + threadIdx.y;
    int j = blockIdx.x * blockDim.x + threadIdx.x;
    int k;
    int sum = 0;
    int dct_i = threadIdx.y;
    int dct_j = threadIdx.x;

    __shared__ int temp[8][8];

    temp[dct_i][dct_j] = src[i*width+j];
    __syncthreads();
    sum = 0;

    for (k=0; k<8; k++) {
        sum += temp[dct_i][k] * dDCT_bases[dct_j*8+k]; 
    }
    __syncthreads();
    temp[dct_i][dct_j] = sum >> 13;
    __syncthreads();

    sum = 0; 

    for (k = 0; k < 8; k++) {
        sum += dDCT_bases[dct_i*8+k] * temp[k][dct_j];
    }
    __syncthreads();

    src[i*width+j] = sum >> 13;

}

void myDCT2D(int *src, int width, int height) {

    int bi, bj;
    int i, j, k;
    int sum = 0;
    int temp[64];

    for (bi = 0; bi < width / 8; bi++) {
        for (bj = 0; bj < height / 8; bj++) {
            for (i=0; i<8; i++) {
                for (j=0; j<8; j++) {
                    for (k = 0; k < 8; k++) {
                        sum +=  src[i*8+k] * DCT_bases[j*8+k];
                    }
                    temp[i*8+j] = sum >> 13;
                    sum = 0;
                }
            }

            for (i=0; i<8; i++) {
                for (j=0; j<8; j++) {
                    for (k=0; k < 8; k++) {
                        sum += DCT_bases[i*8+k] * temp[k*8+j];
                    }
                    src[i*8+j] = sum >> 13;
                    sum = 0;
                }
            }

            src += 64; 
        }
    }       
}

int main (int argc, char *argv[]) 
{
    int *matrix;
    int *m0;
    int i, j;
    int *d_m;
    int *m1;

    FILE* fp;

    int width = 8;
    int height = 8;

    if (argc > 1) { 
        width = atoi(argv[1]);
        height = atoi(argv[2]);
    }


    if (width % 8 || height % 8) {
        printf("Width and Height has to be multiple of 8!\n"); 
        getchar();
        return 0;
    }

    matrix = (int *) malloc(sizeof(int) * width * height); 
    m0 = (int *) malloc(sizeof(int) * width * height);
    m1 = (int *) malloc(sizeof(int) * width * height);
    fp = fopen("cuda_test.txt", "w");

    for (i=0; i< height; i++) {
        for (j = 0; j < width; j++) {
            matrix[i*width+j] = rand()% 256;
            m0[i*width+j] = matrix[i*width+j];
            m1[i*width+j] = matrix[i*width+j];
            fprintf(fp,"%d ", m0[i*width+j]);
        }
        fprintf(fp,"\n");
    }
    fprintf(fp, "\n");

    cudaMalloc((void**) &d_m, sizeof(int) * width * height);
    cudaMemcpy(d_m, m0, sizeof(int) * width * height, cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(dDCT_bases, DCT_bases, sizeof(DCT_bases));

    //    printf("%s\n", cudaGetErrorString(cudaGetLastError()));
    dim3 dimGrid(width / 8, height / 8);    
    dim3 dimBlock(8,8);

    cudaDCT2D<<<dimGrid,dimBlock>>> (d_m, width);

    cudaMemcpy(m0, d_m, sizeof(int) * width * height, cudaMemcpyDeviceToHost);

    for (i=0; i< height; i++) {
        for (j = 0; j < width; j++) {
            fprintf(fp,"%d ", m0[i*width+j]);
        }
        fprintf(fp,"\n");
    }

    fprintf(fp, "\n"); 
    myDCT2D(m1, width, height);

    for (i=0; i< height; i++) {
        for (j = 0; j < width; j++) {
            fprintf(fp,"%d ", m1[i*width+j]);
        }
        fprintf(fp,"\n");
    }

    fprintf(fp,"\n");
    free(matrix);
    free(m0);
    free(m1);
    cudaFree(d_m);

    return 0;
}
4

1 回答 1

1

我自己找到答案。实际上 cuda 程序没有任何问题，但我以不同的方式解释矩阵。在 CUDA 中，我使用二维块结构，因此 cuda 会以这种方式解释 16x16 矩阵： [ M1_8x8 M2_8x8 M3_8x8 M4_8x8] 但在我的 C 测试程序中，我假设前 8x8 数字在第一个矩阵，所以它变成这样：[M1 16x4 M2 16x4 M3 16x4 M4 16x4]

所以矩阵不同！这就是为什么结果不一样的原因！我认为这只会发生在像我这样的初学者身上...... :(

于 2012-12-11T20:33:56.040 回答