2

私はCUDAプログラミングを学び始めたばかりで、リダクションについて混乱しています。

グローバル メモリは共有メモリに比べてアクセス遅延が大きいことはわかっていますが、共有メモリと同様の動作を (少なくとも) シミュレートするためにグローバル メモリを使用できますか?

たとえば、長さが正確にある大きな配列の要素を合計したいBLOCK_SIZE * THREAD_SIZE(グリッドとブロックの次元はどちらも のべき乗です2)。以下のコードを使用しようとしました:

    __global__ void parallelSum(unsigned int* array) {

    unsigned int totalThreadsNum = gridDim.x * blockDim.x;
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;

    int i = totalThreadsNum / 2;
    while (i != 0) {
            if (idx < i) {
                array[idx] += array[idx + i];
        }
        __syncthreads();
        i /= 2;
    }
}

このコードの結果と、ホスト上で連続して生成された結果を比較したところ、奇妙なことに、結果が同じ場合もあれば、明らかに異なる場合もあります。ここでグローバルメモリを使用することに関連する理由はありますか?

4

2 に答える 2

5

トムはすでにこの質問に答えています。彼の回答では、ThrustまたはCUBを使用して CUDA のリダクションを実行することを推奨しています。

ここでは、両方のライブラリを使用してリダクションを実行する方法について、完全に機能する例を提供しています。

#define CUB_STDERR

#include <stdio.h>

#include <thrust/device_ptr.h>
#include <thrust/reduce.h>
#include <thrust/execution_policy.h>

#include <cub/device/device_reduce.cuh>

#include "TimingGPU.cuh"
#include "Utilities.cuh"

using namespace cub;

/********/
/* MAIN */
/********/
int main() {

    const int N = 8388608;

    gpuErrchk(cudaFree(0));

    float *h_data       = (float *)malloc(N * sizeof(float));
    float h_result = 0.f;

    for (int i=0; i<N; i++) {
        h_data[i] = 3.f;
        h_result = h_result + h_data[i];
    }

    TimingGPU timerGPU;

    float *d_data;          gpuErrchk(cudaMalloc((void**)&d_data, N * sizeof(float)));
    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    /**********/
    /* THRUST */
    /**********/
    timerGPU.StartCounter();
    thrust::device_ptr<float> wrapped_ptr = thrust::device_pointer_cast(d_data);
    float h_result1 = thrust::reduce(wrapped_ptr, wrapped_ptr + N);
    printf("Timing for Thrust = %f\n", timerGPU.GetCounter());

    /*******/
    /* CUB */
    /*******/
    timerGPU.StartCounter();
    float           *h_result2 = (float *)malloc(sizeof(float));
    float           *d_result2; gpuErrchk(cudaMalloc((void**)&d_result2, sizeof(float)));
    void            *d_temp_storage = NULL;
    size_t          temp_storage_bytes = 0;

    DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N);
    gpuErrchk(cudaMalloc((void**)&d_temp_storage, temp_storage_bytes));
    DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N);

    gpuErrchk(cudaMemcpy(h_result2, d_result2, sizeof(float), cudaMemcpyDeviceToHost));

    printf("Timing for CUB = %f\n", timerGPU.GetCounter());

    printf("Results:\n");
    printf("Exact: %f\n", h_result);
    printf("Thrust: %f\n", h_result1);
    printf("CUB: %f\n", h_result2[0]);

}

CUB はアルゴリズムの正確な選択や同時実行性の程度などのパフォーマンスに重要な詳細をユーザーの手に委ねているため、CUB は異なる根本的な哲学により、Thrust よりもいくらか高速になる可能性があることに注意してください。このようにして、これらのパラメータを調整して、特定のアーキテクチャとアプリケーションのパフォーマンスを最大化できます。

配列のユークリッド ノルムを計算するための比較は、CUB in Action – CUB テンプレート ライブラリを使用した簡単な例で報告されています。

于 2015-08-08T21:37:10.910 に答える
4

最善の策は、CUDA サンプルのリダクションの例から始めることです。スキャンの例 は、スループット アーキテクチャでの並列計算の原則を学習するのにも適しています。

つまり、実際にコードでリダクション演算子を使用したい場合は、 Thrust (ホストからの呼び出し、クロスプラットフォーム) とCUB (CUDA GPU 固有) を確認する必要があります。

特定の質問を見るには：

  • 削減のためにグローバル メモリを使用できない理由はありません。ツールキットのサンプル コードはさまざまなレベルの最適化を説明していますが、いずれの場合もデータはグローバル メモリから開始されます。
  • あなたのコードは非効率的です (作業効率の詳細については、ツールキットの例を参照してください!)。
  • あなたのコードは、適切な同期なしで異なるブロック内のスレッド間で通信しようとしています。__syncthreads()特定のブロック内のスレッドのみを同期し、異なるブロック間では同期しません (GPU をオーバーサブスクライブする傾向があるため、少なくとも一般的には不可能です)。つまり、特定の時間にすべてのブロックが実行されるわけではありません)。

最後のポイントが最も重要です。ブロック X のスレッドがブロック Y によって書き込まれたデータを読み取りたい場合、これを 2 つのカーネル起動にまたがって分割する必要があります。これが、典型的な並列削減がマルチフェーズ アプローチを採用する理由です。つまり、ブロック内でバッチを削減し、次にバッチ間で削減します。 .

于 2013-12-02T10:44:32.513 に答える