メタルカーネルは、可変データハザードでインデックスされている場合、MTLBufferに乱数を書き込みますか？

Question

私はヒストグラムを計算する必要があるGPU上で高ダイナミックレンジアルゴリズムを実装しようとしています。金属コードは、これまでに次のようになります。

kernel void 
hist(texture2d_array<half, access::read> inArray [[texture(0)]], 
     device float *t [[buffer(0)]], // ignore this 
     volatile device uint *histogram [[buffer(1)]], 
     uint2 gid [[thread_position_in_grid]]){ 

int4 Y_ldr; 
uint redChannel; 

for(uint i = 0; i < inArray.get_array_size(); i++){ 
    Y_ldr = int4(inArray.read(gid, i, 0) * 255); 
    redChannel = Y_ldr.r; 
    histogram[redChannel]++; 
} 

}

カーネルは、膨大な数とヒストグラム（256項目）の半分を埋め、残りの半分は（初期値）はnullです。私は

histogram[0] = 1; // just a number 
histogram[0] = redChannel; // OR this 

を書くとき 代わりに、私は両方のケースでは位置0で正しい番号を取得します。 atomic_uintを使用してもスレッドグループバリアは役に立ちません。

histogram[0]++; 

をしようとすると、

1. は、私は無理番号
を得るというトラブルを引き起こしているものを、そうのような12000 すなわち、金属は、自動的にデータハザードを処理しないことが明らかになったが、数字が小さいです
2. アレイのちょうど半分が欠落していますか？あなたが知る必要がある場合

、パイプライン状態の設定方法は、こちらを参照してください。

var threadGroupCount = MTLSizeMake(8, 8, 1) 
var threadgroups = MTLSizeMake(pictures!.width/threadGroupCount.width, pictures!.height/threadGroupCount.height, 1) 
computeCommandEncoder.setComputePipelineState(hist!) 
computeCommandEncoder.setTexture(pictures, atIndex: 0) 
computeCommandEncoder.setBuffer(exposure_times, offset: 0, atIndex: 0) 
computeCommandEncoder.setBuffer(histogram, offset: 0, atIndex: 1) // <-- this is important!!!CommandEncoder.dispatchThreadgroups(threadgroups, threadsPerThreadgroup: threadGroupCount) 

Answer 1

イエス-.- のiPadのCPUは64ビット数（8バイト）としてUINTを解釈します。 GPUでは、Intは32ビット（4バイト）を意味します。ポインタが[1]の位置に設定されると、CPUはこの位置を配列エントリ[0]の上位4バイトと解釈します。それは非常に大きな数字になります。 この問題を解決するには、ヒストグラムをCPU側で[UInt32]に設定する必要があります。