OpenCLでローカルメモリを宣言するには？

Question

私はこれはしかし、パラメータlengthいるので、コンパイルすることを拒否1000000×100の2次元グローバルワークサイズと×100 1

__kernel void myKernel(
     const int length, 
     const int height, 
     and a bunch of other parameters) { 

    //declare some local arrays to be shared by all 100 work item in this group 
    __local float LP [length]; 
    __local float LT [height]; 
    __local int bitErrors = 0; 
    __local bool failed = false; 

    //here come my actual computations which utilize the space in LP and LT 
} 

のローカルワークサイズを下回るのOpenCLカーネルを実行していることだし、 heightはコンパイル時には分かりません。しかし、これをどうやって正しく行うかは、私には全く分かりません。私はmemallocでポインタを使うべきですか？ワークグループ全体で1回だけメモリが割り当てられ、ワークアイテムごとに1回は割り当てられません。

私が必要とするのは、ワークグループ全体で共有される2つの浮動小数点配列、1 int、1ブール値です（100個すべての作業項目）。

kernel void myKernel(const int length, const int height, local float* LP, 
        local float* LT, a bunch of other parameters) 

あなたは、その後のvalueでkernelargumentを設定します。しかし、私はそれはあなたがあなたのカーネルの引数としてローカル配列を渡すことができ、比較的簡単です

Answer 1

...正しくこれを行う任意の方法を見つけることができませんNULLとsizeは、引数に割り当てるサイズ（バイト単位）と同じです。したがって、それは次のようになります。

clSetKernelArg(kernel, 2, length * sizeof(cl_float), NULL); 
clSetKernelArg(kernel, 2, height* sizeof(cl_float), NULL); 

ローカルメモリは常に（プライベートではなく）ワークグループで共有するので、私はboolとintは問題ないはずと思いますが、ない場合は、常にあまりにも引数としてそれらを渡すことができています。

実際にあなたの問題に関連しているわけではありません（そして、私はあなたがこれを実行する予定のハードウェアを知らないので、必ずしも関連しているわけではありません）。しかし、少なくともgpusは特定の力の倍数ではない2つ（私はそれがNVIDIAでは32、amdでは64）と言います。これは、おそらく最後の28が基本的に無駄になる128項目のワークグループを作成することを意味します。ですから、GPU上のOpenCLを実行している場合は、直接側の注意点としてはサイズ128（および適切グローバルワークサイズを変更する）

のワークグループを使用する場合、それはパフォーマンスを助けるかもしれない：誰もがkernel, local and global用アンダースコアバリアントを使用して、なぜ私は理解していませんでした、私にとってもっと醜いようです。

Answer 2

ローカルメモリをすべてカーネルの外に割り当てる必要はありません。特に、配列ではなく単純な変数の場合は特にそうです。

コードがコンパイルできないのは、OpenCLがローカルメモリの初期化をサポートしていないためです。これは文書（https://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/local.html）で指定されています。これは、CUDAにも（Is there a way of setting default value for shared memory array?）現実的ではありません

---

PS：グリズリーからの答えは十分であると私はコメントとして投稿することができた場合、それは良いでしょうが、私は評判のポリシーで制限しています。ごめんなさい。

Answer 3

また、このようなあなたの配列を宣言できます。

__local float LP[LENGTH]; 

そして、あなたのカーネルのコンパイルに定義するようLENGTHを渡します。

clBuildProgram(program, 0, NULL, "-DLENGTH=128", NULL, NULL);