0

2 つの行列 NxN サイズを乗算する CUDA カーネルを作成したいと考えています。なんとかできたのですが、スレッド連携なしで...今度はスレッド連携でやりたいと思い、SDKで提供されているコードに従いました。しかし、何らかの理由で、カーネルは異なる結果を返します。.cu ファイルは次のとおりです。

#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>
#include<cuda_runtime_api.h>
#include<device_functions.h>

static void HandleError(cudaError_t err, const char *file, int line)
{
    if(err!=cudaSuccess){
    printf("%s in %s file at line %s\n", cudaGetErrorString(err), file, line);
    exit(EXIT_FAILURE);
    }
}

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

#ifndef _MATRIXMUL_KERNEL_H_
#define _MATRIXMUL_KERNEL_H_

#define ORDER 4

__global__ void matrixMul( int* A, int* B, int* C, int wA, int wB)
{
    int bx = blockIdx.x;
        int by = blockIdx.y;

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


    int aBegin = wA * ORDER * by;

    int aEnd   = aBegin + wA - 1;

    int aStep  = ORDER;

    int bBegin = ORDER * bx;

    int bStep  = ORDER * wB;

    int Csub=0;

    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    {
        __shared__ int As[ORDER][ORDER];

        __shared__ int Bs[ORDER][ORDER];

        As[ty][tx] = A[a + wA * ty + tx];
        Bs[ty][tx] = B[b + wB * ty + tx];

        __syncthreads();

        #pragma unroll

        for (int k = 0; k < ORDER; ++k)
            Csub += As[ty][k] * Bs[k][tx];

        __syncthreads();
    }

    int c = wB * ORDER * by + ORDER * bx;
    C[c + wB * ty + tx] = Csub;
}

#endif

int main()
{
    int *a=(int*)malloc(ORDER*ORDER*sizeof(int));
    int *b=(int*)malloc(ORDER*ORDER*sizeof(int));
    int *c=(int*)malloc(ORDER*ORDER*sizeof(int));

    int *dev_a, *dev_b, *dev_c;

    HANDLE_ERROR(cudaMalloc((void**)&dev_a, ORDER*ORDER*sizeof(int*)));
    HANDLE_ERROR(cudaMalloc((void**)&dev_b, ORDER*ORDER*sizeof(int*)));
    HANDLE_ERROR(cudaMalloc((void**)&dev_c, ORDER*ORDER*sizeof(int*)));

    for(int i=0; i<ORDER*ORDER; i++)
    {
        a[i]=1;
        b[i]=2;
    }

    HANDLE_ERROR(cudaMemcpy(dev_a, a, ORDER*ORDER*sizeof(int), cudaMemcpyHostToDevice));
    HANDLE_ERROR(cudaMemcpy(dev_b, b, ORDER*ORDER*sizeof(int), cudaMemcpyHostToDevice));

    matrixMul<<<ORDER, ORDER>>>(dev_a, dev_b, dev_c, ORDER, ORDER);

    HANDLE_ERROR(cudaMemcpy(c, dev_c, ORDER*ORDER*sizeof(int), cudaMemcpyDeviceToHost));

    for(int i=0; i<ORDER*ORDER; i++)
    {
        if((i%ORDER)==0)
            printf("\n\n");
        printf("%d\t", a[i]);
    }

    for(int i=0; i<ORDER*ORDER; i++)
    {
        if((i%ORDER)==0)
            printf("\n\n");
        printf("%d\t", b[i]);
    }

    for(int i=0; i<ORDER*ORDER; i++)
    {
        if((i%ORDER)==0)
            printf("\n\n");
        printf("%d\t", c[i]);
    }

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

はい、「本当の」質問がないことは知っています...しかし、誰かが私を正しい方向に向けることができれば、私は感謝します。ありがとうございました！

さらにコード例が必要な場合はお知らせください。質問を編集します。

編集＃1：言及するのを忘れていました... Visual Studi 2010でnvccを実装できなかったため、デバッガーを使用できません。それについて何か提案はありますか？

編集 #2 : CUDA カーネルとメインの両方が表示されるように質問を更新しました。

4

1 に答える 1

1

thread-geometry が の場合、カーネルは正しいようですBLOCKSIZE x BLOCKSIZE。そうですか？

それが問題でない場合:

スレッド同期なしで動作すると言ったので、おそらくメモリ割り当てが正しいでしょう。

4x4 のスレッド ジオメトリと次の 2 つのマトリックスでテストしてみてください。

1 1 1 1    1 0 0 0
2 2 2 2    0 1 0 0
3 3 3 3    0 0 1 0
5 5 5 5    0 0 0 1

出力は、何が問題なのかについてのヒントを提供するはずです。

于 2012-04-20T06:39:20.057 に答える