오버 헤드를 실제 스레드 작업보다 비싸게 보낼 수 있습니까?

Question

최근 하드 최적화의 가능성에 대해 생각하고 있습니다. 나는 당신이 때로는 무언가를 얻기 위해 3 번의 반복에서 루프를 하드 코딩 할 때 이러한 종류의 최적화를 의미합니다.

한 가지 생각이 내 마음에 들었습니다. 1024 요소의 버퍼가 있다고 가정 해보십시오. 우리는 그것의 모든 단일 요소에 2를 곱하려고합니다. 우리는 버퍼, outBuffer, 크기 (우리가 범위를 벗어 났는지 확인하기 위해)와 [[thread_position_in_grid]]을 전달하는 간단한 커널을 만듭니다. 그런 다음 간단한 곱셈을하고 그 수를 다른 버퍼에 씁니다.

이 같은 조금 보이는 것 :

kernel void multiplyBy2(constant float* in [[buffer(0)]], 
          device float* out [[buffer(1)]], 
          constant Uniforms& uniforms [[buffer(2)]], 
          uint gid [[thread_position_in_grid]]) 
{ 

    if (gid >= uniforms.buffer_size) { return; } 
    out[gid] = in[gid] * 2.0; 
} 

그것에 의해 생산되는 여전히 실제 스레드 작업의 가치 오버 헤드가 파견 내가 만약에 대한 걱정 것은 무엇입니까?

그것은

out[gid * 4 + 0] = in[gid + 0] * 2.0; 
    out[gid * 4 + 1] = in[gid + 1] * 2.0; 
    out[gid * 4 + 2] = in[gid + 2] * 2.0; 
    out[gid * 4 + 3] = in[gid + 3] * 2.0; 

그래서 스레드가 더 이상 조금 작업 할 수와 같은 일을하는, 예를 들어, 4 배 적은 스레드를 파견하는 것이 더 효과적 일 것인가? 또는 가능한 한 실을 얇게하는 것이 더 낫습니다.

Answer 1

예, 이는 인위적인 예제뿐만 아니라 일부 실제 시나리오에서도 마찬가지입니다.

매우 간단한 커널의 경우 디스패치 오버 헤드로 인해 수행 할 작업이 줄어들 수 있지만 성능에 더 큰 영향을 줄 수있는 또 다른 요소가 있습니다. 가져온 데이터와 중간 결과 공유입니다.

예를 들어, 입력 텍스쳐에서 픽셀의 3x3 근방을 읽고 출력 텍스쳐에 평균을 쓰는 경우, 다음과 같이 연산하여 페치 된 텍스처 데이터와 인접한 픽셀 사이의 부분 합계를 공유 할 수 있습니다. 커널 함수의 픽셀이 두 개 이상이고 디스패치하는 총 스레드 수를 줄입니다.

아마도 이것은 호기심을 나타냅니다. 실용적인 응용 프로그램의 경우, Scott Hunter는 최적화 전후에 모든 대상 장치에서 프로파일 링해야합니다.