OpenCL：可変長の配列を出力する

Question

私たちはGPGPUコースの課題に取り組んでいます。私たちはアルゴリズムを選んでCPUに実装し、OpenCLに変換しています。

私たちが選択したアルゴリズムは、モデルを三角形のセットとして読み込み、それらをボクセルにラスタライズします。ボクセルはポイントデータのVBOとして定義されます。その後、ジオメトリシェーダを使用して、これらの点を三角形のボクセルに変換します。

したがって、私たちのOpenCLプログラムは、三角形のリストを取り、点の変数リストを出力する必要があります。

可変長配列を出力することは問題であるようです。

解決策は、カウンタをアトミックにインクリメントし、そのカウンタを出力配列のインデックスと配列の最終サイズの両方として使用することです。例外を除き...両方のGPUは、アトミック操作の拡張をサポートしていません。

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable 

#define POS1  i0 * 3 + 0 
#define POS2  i0 * 3 + 1 
#define POS3  i0 * 3 + 2 

void WritePosition(__global float* OutBuffer, uint inIndex, __global float* inPosition) 
{ 
    OutBuffer[ inIndex * 3 ] = inPosition[0]; 
    OutBuffer[ inIndex * 3 + 1] = inPosition[1]; 
    OutBuffer[ inIndex * 3 + 2] = inPosition[2]; 
} 

__kernel void Voxelize( 
    __global float* outPointcloudBuffer, 
    __global float* inTriangleBuffer, 
    __global uint* inoutIndex 
) 
{ 
    size_t i0 = get_global_id(0); 
    size_t i1 = get_local_id(0); 

    WritePosition(outPointcloudBuffer, inIndex[0], &inTriangleBuffer[ i0 ]); 

    //atomic_inc(inoutIndex[0]); 
    inoutIndex[0] = max(inoutIndex[0], i0); 
} 

そして、これの出力は非常に奇妙です：

これは、我々がこれまで持っているものです。我々は非常に小さなモデル（12三角形、36ポジション、108フロート）をテストしており、結果は31,63、または95です。常に16の倍数から1を引いたものです。

私たちの可変長出力配列？

ありがとうございます。

Answer 1

Iは次のようにこれを正常に取り組んでいることを推測する：

• 最初のパスは：プリミティブscan（並列プレフィックス合計）を使用してGPU上でアレイの必要なサイズを計算します。上のリンクには、Appleの実装例が含まれています。
• スキャンアルゴリズムの結果を使用して、ホスト側で必要なリソースを割り当てます。スキャンアルゴリズムの結果は、個々の作業項目の結果のインデックスヒントとしてよく使用されることに注意してください。
• 2回目のパス（オプション）：3回目のパスで考慮する必要のある要素に配列を圧縮します。
• 第3パス：宛先インデックスと割り当てられた配列を渡すアルゴリズムを再実行します。

上記の3つのパスがすべて実装されているNVIDIAのOpenCLマーチングキューブimplementationをご覧ください。

ベスト、Christoph