플로트 버퍼

는

문제점 : I는 상수 값으로 MTLBufferFloat의 S 채울 필요플로트 버퍼

는 - 1729.68921 말한다. 나는 또한 가능한 한 빨리해야한다.

따라서 CPU 쪽에서 버퍼를 채우지 않아야합니다 (즉, UnsafeMutablePointer<Float>을 MTLBuffer에서 얻고 직렬 방식으로 지정).

내 접근

이상적으로는 UInt8는 1 바이트 길이이며 Float, 내가 할 수없는 4 바이트 길이 주어진 UInt8 값 (로 버퍼를 채우기 만 할 수 있어요하지만 AFAIK, MTLBlitCommandEncoder.fill()을 사용하십시오 내 Float 상수 값 임의로 지정하십시오.

지금까지 나는 단지이 옵션은 왼쪽 볼 수 있지만 모두 잔인한 것 같다 :

상수 값 가득 다른 버퍼 B을 만들고 MTLBlitCommandEncoder
을 만들 통해 내 버퍼에 그 내용을 복사 버퍼를

질문

채울 것 kernel 함수

이 Float 인 가장 빠른 방법은 이고 상수 값은 무엇입니까?

출처

2016-12-17 sarasvati

GPU에서 상수가 필요한 경우 가장 빠른 옵션은 커널 함수 내에 해당 상수를 만들고 버퍼를 통해 CPU에서 하나를 보내는 대신 로컬에서 사용하는 것입니다. – Marius

@ 마리우스 나는 왜 내가 버퍼를 채워야하는지 명확히해야했다 : 나는 물건을 계산하는 루프를 가지고있다.루프의 시작 부분에서 몇 가지 상수 값으로 초기화 할 버퍼가 필요합니다. 나중에 루프 내에서 버퍼의 여러 위치에 * 다른 값을 할당합니다. 위치는 각 루프마다 다릅니다. – sarasvati

CPU 사용 속도가 느린 것이 확실합니까? 당신은'memset_pattern4()'에 대해 알고 있고 자신의 순진한 구현 이상으로 최적화 될 가능성이 있습니까? –

각 스레드의 여러 버퍼 요소에 쓰는 컴퓨팅 쉐이더를 사용하는 것이 내 실험에서 가장 빠른 방법이었습니다. 이는 하드웨어에 따라 다르므로 앱을 배포 할 것으로 예상되는 모든 기기에서 테스트해야합니다.

kernel void fill_16_unchecked(device float *buffer [[buffer(0)]], 
           constant float &value [[buffer(1)]], 
           uint index   [[thread_position_in_grid]]) 
{ 
    for (int i = 0; i < 16; ++i) { 
     buffer[index * 16 + i] = value; 
    } 
} 

kernel void single_fill_checked(device float *buffer   [[buffer(0)]], 
           constant float &value  [[buffer(1)]], 
           constant uint &buffer_length [[buffer(2)]], 
           uint index     [[thread_position_in_grid]]) 
{ 
    if (index < buffer_length) { 
     buffer[index] = value; 
    } 
}

당신 경우 : 버퍼의 길이에 대해 검사 후 단일 배열 요소를 설정하는 배열 범위에 대해 검사하지 않고 16 개 개의 연속 배열 요소를 채우고 하나 및 하나

나는 두 연산 쉐이더 쓴 당신의 버퍼 카운트가 항상 스레드 실행 너비에 루프에서 설정 한 요소의 수를 곱한 값의 배수가된다는 것을 알고 있다면, 첫 번째 함수를 사용할 수 있습니다. 두 번째 함수는 그렇지 않으면 버퍼를 오버런시키는 그리드를 디스패치 할 수있는 경우의 대체 기능입니다.

이러한 함수에서 내장 된 두 개의 파이프 라인이 있으면, 다음과 같이 계산 명령 한 쌍의 작업을 전달할 수 있습니다

NSInteger executionWidth = [unchecked16Pipeline threadExecutionWidth]; 
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder]; 
[computeEncoder setBuffer:buffer offset:0 atIndex:0]; 
[computeEncoder setBytes:&value length:sizeof(float) atIndex:1]; 
if (bufferCount/(executionWidth * 16) != 0) { 
    [computeEncoder setComputePipelineState:unchecked16Pipeline]; 
    [computeEncoder dispatchThreadgroups:MTLSizeMake(bufferCount/(executionWidth * 16), 1, 1) 
        threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)]; 
} 
if (bufferCount % (executionWidth * 16) != 0) { 
    int remainder = bufferCount % (executionWidth * 16); 
    [computeEncoder setComputePipelineState:checkedSinglePipeline]; 
    [computeEncoder setBytes:&bufferCount length:sizeof(bufferCount) atIndex:2]; 
    [computeEncoder dispatchThreadgroups:MTLSizeMake((remainder/executionWidth) + 1, 1, 1) 
        threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)]; 
} 
[computeEncoder endEncoding];

이런 식으로 일을 반드시 빨리 초과 할 수 없습니다 것이다하는 것으로 스레드 당 하나의 요소 만 쓰는 단순한 접근 방식. 테스트에서 A8에서 40 % 빨라졌으며 A10에서는 대략 같았고 A9에서는 2-3 배 느려졌습니다. 항상 자신의 작업 부하로 테스트하십시오.

출처

2016-12-18 21:36:26 warrenm

플로트 버퍼

답변

관련 문제