1

私はCUDAを学び始めたばかりで、非常に基本的なことしか知りません。行列乗算法を使用して8x8DCTを実行する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では2Dブロック構造を使用しているため、16x16マトリックスはcudaによって次のように解釈されます。[M1_8x8 M2_8x8 M3_8x8 M4_8x8]しかし、私のCテストプログラムでは、最初の8x8の数値が最初のマトリックスなので、次のようになります。[M1 16x4 M2 16x4 M3 16x4 M4 16x4]

したがって、マトリックスは異なります！そのため、結果は同じではありません。私のような初心者にのみ起こると思います...:(

于 2012-12-11T20:33:56.040 に答える