CL_DEVICE_MAX_WORK_GROUP_SIZEはなぜですか？

Question

私はGPUなどのOpenCLデバイスのアーキテクチャを理解しようとしていますが、ローカルワークグループ内の作業項目の数、つまり定数CL_DEVICE_MAX_WORK_GROUP_SIZEに明示的な境界がある理由がわかりません。

これはコンパイラによって処理されるべきであると思われます。つまり、物理的な最大値が100でありながらカーネルがローカルワークグループサイズ500で実行され、カーネルが

__kernel void test(float* input) { 
    i = get_global_id(0); 
    someCode(i); 
    barrier(); 
    moreCode(i); 
    barrier(); 
    finalCode(i); 
} 

それは、このカーネル上のワークグループサイズ100で実行に自動的に変換することができます：このよう

__kernel void test(float* input) { 
    i = get_global_id(0); 
    someCode(5*i); 
    someCode(5*i+1); 
    someCode(5*i+2); 
    someCode(5*i+3); 
    someCode(5*i+4); 
    barrier(); 
    moreCode(5*i); 
    moreCode(5*i+1); 
    moreCode(5*i+2); 
    moreCode(5*i+3); 
    moreCode(5*i+4); 
    barrier(); 
    finalCode(5*i); 
    finalCode(5*i+1); 
    finalCode(5*i+2); 
    finalCode(5*i+3); 
    finalCode(5*i+4); 
} 

しかし、これはデフォルトで実行されていないようです。何故なの？このプロセスを自動化する方法はありますか（それのためにプリコンパイラを書くこと以外にも）？それとも、私のメソッドが特定の例で失敗するようにする本質的な問題がありますか？

Answer 1

私はCL_DEVICE_MAX_WORK_GROUP_SIZEの起源は、基本的なハードウェアの実装にあると思います。

複数のスレッドが計算ユニットで同時に実行されており、すべてのスレッドが状態を保持する必要があります（呼び出し、jmpなど）。ほとんどの実装ではこのためにスタックが使用されています.Evergreenファミリを見ると、使用可能なスタックエントリの数のハードウェア制限があります（スタックエントリごとにサブエントリがあります）。これは本質的に、各コンピューティングユニットが同時に処理できるスレッドの数を制限します。

コンパイラはこれを可能にするためにこれを行うことができます。それはうまくいくかもしれませんが、カーネルを再コンパイルすることを意味することを理解しています。これは常に可能ではありません。開発者がコンパイルされたカーネルを各プラットフォーム用にバイナリ形式でダンプし、そのソフトウェアと共に「あまりオープンソースでない」理由で出荷する状況を想像することができます。

Answer 2

これらの定数は、コンパイル時に適切なワークグループサイズを決定するために、コンパイラによってデバイスから照会されます（コンパイルはもちろんカーネルのコンパイルを指します）。私はあなたが間違っているかもしれないかもしれませんが、あなたは自分でこれらの値を設定することを考えているようですが、それはそうではありません。

実行するハードウェアに応じて準備するシステム機能を照会する責任は、コード内にあります。