금속 격자의 모든 스레드 동기화

Question

Metal에서 n 크기 벡터에 대한 표준 또는 제곱 길이 함수를 작성하려고합니다. 이를 위해 모든 스레드가 각 요소를 정사각형으로 만들도록 계획 한 다음 하나의 스레드를 선택하여 모든 요소를 ​​합산했습니다.

#include <metal_stdlib> 
#include <metal_compute> 
using namespace metal; 

kernel void length_squared(const device float *x [[ buffer(0) ]], 
          device float *s [[ buffer(1) ]], 
          device float *out [[ buffer(2) ]], 
          uint gid [[ thread_position_in_grid ]], 
          uint numElements [[ threads_per_grid ]]) 
{ 
    s[gid] = x[gid];// * x[gid]; 
    simdgroup_barrier(mem_flags::mem_none); 
    if(gid == 0){ 
     for(uint i = 0; i < numElements; i++){ 
      *out += s[i]; 
     } 
    } 
} 

불행하게도,이 코드는 컴파일되지 않는다 "선언되지 않은 식별자 simdgroup_barrier의 사용"에 대한 :

여기에 내 현재 커널입니다. 이 방법은 Metal Shading Language Specification에 설명되어 있습니다.

누구에게이 문제가 발생 했습니까? 또는 그리드에서 모든 스레드를 동기화하는 방법을 알고 있습니까? threadgroup_barrier이 나를 위해 전체 동기화를 수행하지 않습니다.

이 문제는 제대로 수행되지 않습니까? 이 작업을 동기화하는 가장 좋은 방법은 무엇입니까?

Answer 1

SIMD 그룹이 스레드 그룹보다 작으므로 동기화가 작동하지 않습니다.

대신에 parallel reduction을 사용하여 병렬로 값을 합산 할 수 있습니다. Here은 내가 발견 한 금속 코드입니다.

하나의 스레드가 모든 합계를 수행하는 것에 신경 쓰지 않는다면 합계를 수행하기 위해 하나의 스레드만으로 별도의 커널을 실행할 수 있습니다. 물론 이것은 매우 느릴 수 있습니다.