Расчет среднего значения в металлическом ядре

Question

Кто-нибудь знает правильный способ вычисления среднего значения буфера со случайными числами с плавающей точкой в ​​металлическом ядре?

Диспетчерская работа по командной вычислений кодера:

threadsPerGroup = MTLSizeMake(1, 1, inputTexture.arrayLength); 
numThreadGroups = MTLSizeMake(1, 1, inputTexture.arrayLength/threadsPerGroup.depth); 

[commandEncoder dispatchThreadgroups:numThreadGroups 
       threadsPerThreadgroup:threadsPerGroup]; 

Код ядра:

kernel void mean(texture2d_array<float, access::read> inTex [[ texture(0) ]], 
      device float *means       [[ buffer(1) ]], 
      uint3 id          [[ thread_position_in_grid ]]) { 

if (id.x == 0 && id.y == 0) { 
    float mean = 0.0; 
    for (uint i = 0; i < inTex.get_width(); ++i) { 
     for (uint j = 0; j < inTex.get_height(); ++j) { 
       mean += inTex.read(uint2(i, j), id.z)[0]; 
     } 
    } 

    float textureArea = inTex.get_width() * inTex.get_height(); 
    mean /= textureArea; 
    out[id.z] = mean; 
} 

}

Буфер представлен в текстуру texture2d_array типа с форматом R32Float пикселей.

Answer 1

Если вы можете использовать массив uint (вместо float) в качестве источника данных, я бы предложил использовать «функции Atomic Fetch и Modify» (как описано в языке металлического затенения spec) для атомарного ввода в буфер ,

Вот пример функции ядра, которая принимает входной буфер (данные: массив Float) и записывает сумму буфера в атомный буфер (сумма, указатель на UINT):

kernel void sum(device uint *data [[ buffer(0) ]], 
       volatile device atomic_uint *sum [[ buffer(1) ]], 
       uint gid [[ thread_position_in_grid ]]) 
{ 
    atomic_fetch_add_explicit(sum, data[gid], memory_order_relaxed); 
} 

В вашем быстром файле, вы должны установить буфера:

... 
let data: [UInt] = [1, 2, 3, 4] 
let dataBuffer = device.makeBuffer(bytes: &data, length: (data.count * MemoryLayout<UInt>.size), options: []) 
commandEncoder.setBuffer(dataBuffer, offset: 0, at: 0) 

var sum:UInt = 0 
let sumBuffer = device!.makeBuffer(bytes: &sum, length: MemoryLayout<UInt>.size, options: []) 
commandEncoder.setBuffer(sumBuffer, offset: 0, at: 1) 
commandEncoder.endEncoding() 

COMMIT, подождать и затем извлечь данные из GPU:

commandBuffer.commit() 
commandBuffer.waitUntilCompleted() 

let nsData = NSData(bytesNoCopy: sumBuffer.contents(), 
         length: sumBuffer.length, 
         freeWhenDone: false) 
nsData.getBytes(&sum, length:sumBuffer.length) 

let mean = Float(sum/data.count) 
print(mean) 

В качестве альтернативы, если ваш исходный источник данных должен быть массивом float, вы можете использовать метод ускорения Accelerate, который очень быстро подходит для таких вычислений, с помощью метода vDSP_meanv.

Надеюсь, что это помогло, ура!