1

私は CUDA の学習を開始し、いくつかのデータを GPU にコピーして変更し、元に戻す単純なプログラムを書きたいと思いました。私はすでにグーグルで調べて、自分の間違いを見つけようとしました。問題がカーネルにあることは確かですが、何が問題なのか完全にはわかりません。

これが私のカーネルです：

__global__ void doStuff(float* data, float* result)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[index] = (float) index;
    }
}

そして、これが my の関連部分ですmain:

#include <stdlib.h>
#include <stdio.h>

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int width_in_bytes = 3 * sizeof(float);
    int height = 3;

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16, 16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_blocks>>>(data_array, result_array);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }

    return 0;
}

これを実行する0.000000 1.000000 2.00000と、最初の行が取得され、残りの 2 行がガベージになります。

4

2 に答える 2

3

あなたがcudaを学び始めたばかりなら、2D配列に焦点を当てるかどうかはわかりません。

threads_per_blockまた、変数が定義されているため、コードを質問に手動で入力した場合も興味がありますがthreads_per_blocks、カーネル呼び出しで使用します。

とにかく、コードにはいくつかの問題があります。

  1. 2D 配列を使用する場合、ほとんどの場合、ピッチ パラメータを (何らかの方法で) カーネルに渡す必要があります。 cudaMallocPitch 次の行が適切に整列された境界で始まるように、各行の最後に追加のパディングを使用して配列を割り当てます。これにより、通常、割り当ての粒度は 128 または 256 バイトになります。したがって、最初の行には 3 つの有効なデータ エンティティがあり、その後に 256 バイト (ピッチ変数と同じ) を埋めるのに十分な空きスペースが続きます。そのため、これを考慮してカーネルの呼び出しとカーネル自体を変更する必要があります。
  2. threadIdx.yカーネルは本質的に 1D カーネルです (たとえば、 を理解したり使用したりしません)。したがって、2D グリッドを起動しても意味がありません。この場合は何の問題もありませんが、他のコードでは混乱を招き、面倒な冗長性が生じます。

上記のコメントに基づいて、期待される結果をもたらすいくつかの変更を示す更新されたコードを次に示します。

#include <stdio.h>


__global__ void doStuff(float* data, float* result, size_t dpitch, size_t rpitch, int width)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
    }
}

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int height = 3;
    int width = 3;
    int width_in_bytes = width * sizeof(float);

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_block>>>(data_array, result_array, data_array_pitch, result_array_pitch, width);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }
    return 0;
}

また、この質問を読むと面白いかもしれません。

編集:コメントの質問に答える:

result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
              1               2                      3

ピッチ付き配列への正しい要素インデックスを計算するには、次のことを行う必要があります。

  1. スレッド インデックスから (仮想) 行インデックスを計算します。これを行うには、スレッド インデックスを各 (ピッチのない) 行の幅 (バイト単位ではなく要素単位) で整数除算します。
  2. 行インデックスに各ピッチ行の幅を掛けます。ピッチ付きの各行の幅は、バイト単位のピッチ付きパラメータによって指定されます。このピッチ付きバイトパラメータをピッチ付き要素パラメータに変換するには、各要素のサイズで割ります。次に、手順 1 で計算された行インデックスを数量に掛けることで、正しい行にインデックスを付けました。
  3. スレッド インデックスを幅 (要素単位) で割った剰余 (モジュロ除算) を取得して、スレッド インデックスから (仮想) 列インデックスを計算します。列インデックス (要素単位) を取得したら、それをステップ 2 で計算された正しい行の開始インデックスに追加して、このスレッドが担当する要素を識別します。

上記は、比較的単純な操作のためのかなりの労力です。これは、最初にピッチド アレイではなく、基本的な cuda の概念に焦点を当てることをお勧めする理由の 1 つです。たとえば、ピッチド アレイに取り組む前に、1 次元と 2 次元のスレッド ブロック、および 1 次元と 2 次元のグリッドを処理する方法を考えます。ピッチ付き配列は、場合によっては 2D 配列 (または 3D 配列) にアクセスするための便利なパフォーマンス エンハンサーですが、CUDA で多次元配列を処理する必要はありません。

于 2013-06-03T21:40:03.607 に答える
0

実際には、行を置き換えることによっても実行できます

int width_in_bytes = 3 * sizeof(float);

に：

int width_in_bytes = sizeof(float)*9;

これは、src から dst にコピーするバイト数を cudaMemcpy2D に指示するパラメータであるため、最初のコードでは 3 つの浮動小数点数をコピーするように要求しますが、コピーする配列の長さは 9 であるため、必要な幅は9 つの浮動小数点数。

このソリューションは機能しますが、コードにはまだ非効率な部分があります。たとえば、ブロックの最初の 9 つのスレッドで何かを実行したい場合は、「if」に and(&&) を使用して次の条件を追加する必要があります。

threadIdx.y==0
于 2013-06-03T22:09:50.663 に答える