금속 커널의 평균값 계산

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; 
} 

}

버퍼는 R32Float 픽셀 포맷 texture2d_array 유형의 텍스처로 표현된다.

Answer 1

데이터 소스로 uint (float 대신) 배열을 사용할 수 있다면 금속성 음영 언어 spec에 설명 된대로 "원자 반입 및 수정 기능"을 사용하여 버퍼에 기본적으로 기록하는 것이 좋습니다. . (a 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); 
} 

: 여기

은 (플로트의 배열 데이터) 입력 버퍼를 취하는 커널 함수의 예

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) 

다음 GPU에서 데이터를 가져 오기 한 후,

... 
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() 

커밋 기다린 : 당신의 신속한 파일에서

, 당신은 버퍼를 설정합니다

또는 초기 데이터 소스가 float 배열이어야하는 경우 해당 계산을 위해 매우 빠른 Accelerate 프레임 워크의 vDSP_meanv 메서드를 사용할 수 있습니다.

나는 그것이 도움이 되었으면 좋겠다!