OpenCLのローカルメモリフェンスとグローバルメモリフェンスの違いは何ですか？

Question

私は、例のようにPyOpenCLを使って還元合計を作成しようとしていました：https://dournac.org/info/gpu_sum_reduction。私はすべての値1でベクトルを合計しようとしています。結果は最初の要素で16384になるはずです。 しかし、ちょうどいくつかの点が集まっているようです。ローカルインデックスが必要ですか？競合状態はありますか（結果を同じにしないで2回実行すると）？次のコードで何が間違っていますか？

import numpy as np 
import pyopencl as cl 

def readKernel(kernelFile): 
    with open(kernelFile, 'r') as f: 
     data=f.read() 
    return data 

a_np = np.random.rand(128*128).astype(np.float32) 
a_np=a_np.reshape((128,128)) 
print(a_np.shape) 

device = cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0] 
print(device) 
ctx=cl.Context(devices=[device]) 
#ctx = cl.create_some_context() #ask which context to use 
queue = cl.CommandQueue(ctx) 
mf = cl.mem_flags 

a_g = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=a_np) 

prg = cl.Program(ctx,readKernel("kernel2.cl")).build() 

prg.test(queue, a_np.shape, None, a_g) 

cl.enqueue_copy(queue, a_np, a_g).wait() 
np.savetxt("teste2.txt",a_np,fmt="%i") 

カーネルは次のとおりです。

__kernel void test(__global float *count){ 
    int id = get_global_id(0)+get_global_id(1)*get_global_size(0); 
    int nelements = get_global_size(0)*get_global_size(1); 

    count[id] = 1; 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    for (int stride = nelements/2; stride>0; stride = stride/2){ 
     barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update 
     if (id < stride){ 
      int s1 = count[id]; 
      int s2 = count[id+stride]; 
      count[id] = s1+s2; 
     } 
    } 
    barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update 
} 

Answer 1

問題は、カーネルは、1つのワークグループ内削減を行うために実装され、暗黙的に多くのワークグループがschedulledされていることです。

GPUに応じて、ワークグループごとに異なる最大作業項目数があります。 NVIDIAの場合、1024、AMDとIntel 256（古いGPUのIntelは512でした）。

この例では、GPUのワークグループあたりの最大作業項目が256であると仮定します。この場合、最大2dのworgroupサイズは16x16になる可能性があるため、そのサイズの行列を使用するとカーネルは正しい結果を返します。カーネルのスケジューリング時に元のサイズ128x128を使用し、ローカルサイズを指定しないと、インプリメンテーションで計算されます。グローバルサイズは128x128、ローカルサイズは16x16です。これは8個のグループがスケジュールされていることを意味します。 現在のカーネルでは、各ワークグループは異なるidから計算を開始していますが、インデックスは0になるまで縮小されているため、競合状態が発生し、結果はそれぞれ異なる結果になります。 （16×16）、（16,16）、または任意のあなたの最大の仕事：

1つのワークグループ内のすべてを計算し、グローバル、ローカルの大きさとそれをスケジュールするためにカーネルを書き換え
1. ：

   あなたはこれを修正するための2つのオプションを持っていますワークグループデバイスあたりのアイテム数は

2. です。グローバルなローカルサイズ（128x128）、（16x16）を使用し、各ワークグループはその結果を計算し、最終的な結果を得るために各ワークグループごとに合計する必要があります。

128x128の場合、最初のオプションが優先されます。これは、パフォーマンスが向上し、実装するのがより簡単であるためです。