2

いくつかの例を見て、要素の配列を 1 つの要素に減らしましたが、成功しませんでした。誰かがこれを NVIDIA フォーラムに投稿しました。浮動小数点変数から整数に変更しました。

__kernel void sum(__global const short *A,__global unsigned long  *C,uint size, __local unsigned long *L) {
            unsigned long sum=0;
            for(int i=get_local_id(0);i<size;i+=get_local_size(0))
                    sum+=A[i];
            L[get_local_id(0)]=sum;

            for(uint c=get_local_size(0)/2;c>0;c/=2)
            {
                    barrier(CLK_LOCAL_MEM_FENCE);
                    if(c>get_local_id(0))
                            L[get_local_id(0)]+=L[get_local_id(0)+c];

            }
            if(get_local_id(0)==0)
                    C[0]=L[0];
            barrier(CLK_LOCAL_MEM_FENCE);
}

これは正しく見えますか？3 番目の引数「サイズ」は、ローカル ワーク サイズですか、それともグローバル ワーク サイズですか。

私は自分の主張を次のように設定しました。

clSetKernelArg(ocReduce, 0, sizeof(cl_mem), (void*) &DevA);
clSetKernelArg(ocReduce, 1, sizeof(cl_mem), (void*) &DevC); 
clSetKernelArg(ocReduce, 2, sizeof(uint),   (void*) &size);  
clSetKernelArg(ocReduce, 3, LocalWorkSize * sizeof(unsigned long), NULL);

入力である最初の引数は、その前に起動されたカーネルの出力から保持しようとしています。

clRetainMemObject(DevA);
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocKernel, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
//the device memory object DevA now has the data to be reduced

clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocReduce, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
clEnqueueReadBuffer(hCmdQueue[Plat-1][Dev-1],DevRE, CL_TRUE, 0, sizeof(unsigned long)*512,(void*) RE , 0, NULL, NULL);

今日は、次の cuda リダクションの例を openCL に変換しようと考えています。

__global__ voidreduce1(int*g_idata, int*g_odata){
extern __shared__ intsdata[];

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
__syncthreads();


for(unsigned int s=blockDim.x/2; s>0; s>>=1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}

// write result for this block to global mem
if(tid == 0) g_odata[blockIdx.x] = sdata[0];
}

より最適化されたものがあります (完全にアンロール + スレッドごとに複数の要素)。

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

これはopenCLを使用して可能ですか?

グリズリーは先日こんなアドバイスをくれた

「... n 要素で動作するリダクション カーネルを使用し、それらを n / 16 (またはその他の数値) のようなものに減らします。次に、結果である 1 つの要素になるまで、そのカーネルを繰り返し呼び出します。」

これもやってみたいのですが、どこから始めればいいのかよくわからず、まずは何かを動かしたいです。

4

2 に答える 2

8

最初に指定したリダクション コードは、リダクションに取り組んでいるワークグループが 1 つだけである限り (そうget_global_size(0) == get_local_size(0)) 機能するはずです。その場合、sizeカーネルの引数は要素の数になりますA(これは、グローバルまたはローカルのワークサイズと実際の相関関係はありません)。これは実行可能な解決策ですが、リダクションの実行中にほとんどをアイドル状態にするのは本質的に無駄に思えgpuます。これがまさに、リダクション カーネルを繰り返し呼び出すことを提案した理由です。これは、コードをわずかに変更するだけで可能になります。

__kernel void sum(__global const short *A, __global unsigned long  *C, uint size, __local unsigned long *L) {
        unsigned long sum=0;
        for(int i=get_global_id(0); i < size; i += get_global_size(0))
                sum += A[i];
        L[get_local_id(0)]=sum;

        for(uint c=get_local_size(0)/2;c>0;c/=2)
        {
                barrier(CLK_LOCAL_MEM_FENCE);
                if(c>get_local_id(0))
                        L[get_local_id(0)]+=L[get_local_id(0)+c];

        }
        if(get_local_id(0)==0)
                C[get_group_id(0)]=L[0];
        barrier(CLK_LOCAL_MEM_FENCE);
}

GlobalWorkSizeより小さい then size(たとえば)でこれを呼び出すと、入力がの係数で4減少します。これは反復できます (出力バッファを別の出力バッファでの次の呼び出しの入力として使用することによって)。 、2番目（および後続のすべて）の反復はタイプである必要があるため、実際にはカーネルが必要になりますが、アイデアは得られます.A4*LocalWorkSizesumAglobal const unsigned long*

cuda リダクション サンプルについて: なぜわざわざ変換するのですか? 基本的には、上に投稿した opencl バージョンとまったく同じように機能しますが、イテレーションごとにハードコードされたサイズ ( 2*LocalWorkSizeinsted of size/GlobalWorkSize*LocalWorkSize) だけを縮小することを除きます。

個人的には、削減には実質的に同じアプローチを使用しますが、カーネルを 2 つの部分に分割し、最後の反復ではローカル メモリを使用するパスのみを使用します。

__kernel void reduction_step(__global const unsigned long* A, __global unsigned long  * C, uint size) {
        unsigned long sum=0;
        for(int i=start; i < size; i += stride)
                sum += A[i];
        C[get_global_id(0)]= sum;
}

最後のステップでは、作業グループ内でリダクションを行うフル バージョンが使用されました。もちろん、take の 2 番目のバージョンが必要になります。reduction stepこのglobal const short*コードは、テストされていないコードの適応です (残念ながら、自分のバージョンを投稿することはできません)。wasted workこのアプローチの利点は、ほとんどの作業を実行するカーネルの複雑さが大幅に軽減され、分岐分岐による作業量が少なくなることです。これにより、他のバリアントよりも少し速くなりました。ただし、最新のコンパイラバージョンでも最新のハードウェアでも結果が得られていないため、その点が正しいかどうかはわかりません (ただし、分岐の数が減ったためと思われます)。

リンク先の論文について：openclでサポートされていないテンプレートの使用を除いて、その論文で提案されている最適化をopenclで使用することは確かに可能であるため、ブロックサイズをハードコーディングする必要があります。もちろん、opencl バージョンはすでにカーネルごとに複数の追加を行っており、上記のアプローチに従う場合、ローカル メモリを使用して縮小を展開してもあまりメリットがありません。十分に大きな入力の計算時間全体のかなりの部分。さらに、展開された実装での同期​​の欠如は少し面倒です。その部分に入るすべてのスレッドが同じワープに属しているためにのみ機能します。しかし、これはそうではありません」

于 2012-01-14T20:27:49.513 に答える
1

リダクション カーネルは私の目には正しく見えます。リダクションでは、サイズは入力配列の要素数である必要がありますA。このコードは、スレッドごとの部分合計を に累積しsum、次にローカル メモリ (共有メモリ) 削減を実行して、結果を に格納しCます。Cローカル ワーク グループごとに 1 つの部分合計が得られます。1 つのワーク グループでカーネルをもう一度呼び出して最終的な回答を取得するか、部分的な結果をホストに蓄積します。

于 2012-01-14T20:26:43.730 に答える