3

私は OpenCl の初心者です。

double の 1 次元配列に対してリダクション (合計演算子) を操作する必要があります。

私はネットを徘徊してきましたが、見つけた例はかなり混乱しています。誰でも読みやすい (そしておそらく効率的な) チュートリアルの実装を投稿できますか?

追加情報: - 1 つの GPU デバイスにアクセスできます。- カーネル コードに C を使用しています

4

1 に答える 1

6

あなたの問題には、デバイスのローカルメモリに収まらない60k doubleが含まれていると述べました。ベクトルを 10 ～ 30 程度の値に減らすカーネルをまとめました。これをホスト プログラムで合計できます。私のマシンでは double に問題がありますが、double を有効にし、見つかった場所で「float」を「double」に変更すると、このカーネルは正常に動作するはずです。私が抱えている二重の問題をデバッグし、更新を投稿します。

パラメータ:

  • global float* inVector - 合計する float のソース
  • グローバル float* outVector - 作業グループごとに 1 つずつ、float のリスト
  • const int inVectorSize - inVector が保持する float の総数
  • local float* resultScratch - 各ワーク グループが使用するローカル メモリ。グループ内の作業項目ごとに 1 つのフロートを割り当てる必要があります。予想されるサイズ = sizeof(cl_float)*get_local_size(0)。たとえば、グループごとに 64 個の作業項目を使用する場合、これは 64 float = 256 バイトになります。double に切り替えると 512 バイトになります。openCL 仕様で定義されている LDS の最小サイズは 16kb です。ローカル (NULL) パラメーターを渡す方法の詳細については、この質問を参照してください。

使用法：

  1. 入力バッファーと出力バッファーにメモリを割り当てます。
  2. デバイス上のコンピューティング ユニットごとにワーク グループを作成します。
  3. 最適な作業グループ サイズを決定し、これを使用して「resultScratch」のサイズを計算します。
  4. カーネルを呼び出し、 outVector をホストに読み戻します
  5. outVector のコピーをループして追加し、最終的な合計を取得します。

考えられる最適化:

  1. いつものように、大量のデータを使用してカーネルを呼び出します。データが少なすぎると、転送とセットアップに時間をかける価値がありません。

  2. inVectorSize (およびベクトル) を (ワークグループ サイズ) * (ワーク グループ数) の最大倍数にします。この量のデータだけでカーネルを呼び出します。カーネルはデータを均等に分割します。コールバックを待機している間にホストに残っているデータの合計を計算します (または、CPU デバイス用に同じカーネルを構築し、残りのデータのみを渡します)。上記のステップ 5 で outVector を追加するときは、この合計から始めます。この最適化により、計算全体で作業グループが均等に飽和状態に保たれます。

    __kernel void floatSum(__global float* inVector, __global float* outVector, const int inVectorSize, __local float* resultScratch){
    int gid = get_global_id(0);
    int wid = get_local_id(0);
    int wsize = get_local_size(0);
    int grid = get_group_id(0);
    int grcount = get_num_groups(0);

    int i;
    int workAmount = inVectorSize/grcount;
    int startOffest = workAmount * grid + wid;
    int maxOffest = workAmount * (grid + 1);
    if(maxOffset > inVectorSize){
        maxOffset = inVectorSize;
    }
    resultScratch[wid] = 0.0;
    for(i=startOffest;i<maxOffest;i+=wsize){
            resultScratch[wid] += inVector[i];
    }
    barrier(CLK_LOCAL_MEM_FENCE);

    if(gid == 0){
            for(i=1;i<wsize;i++){
                    resultScratch[0] += resultScratch[i];
            }
            outVector[grid] = resultScratch[0];
    }

    }

また、ダブルを有効にします：

#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#else
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif
#endif

更新: AMD APP KernelAnalyzer が更新 (v12) され、このカーネルの倍精度バージョンが実際には 5870 および 6970 カードにバインドされた ALU であることを示しています。

于 2012-04-16T14:01:18.997 に答える