[Metal] Data Types - SIMD-group Matrix Data Types

마이클 플린이 제안한 플린 분류(Flynn's taxonomy)라는 컴퓨터 아키텍처 분류에 따르면

아키텍쳐는 사용가능한 동시 명령어 수와 데이터 스트림에 따라 나눌 수 있어요.

 

SIMD

 

각각 단일(Single)이냐 다중(Multiple)이냐에 따라서 총 네 가지로 나뉘는데

그 중 SIMD(Single Instruction, Multiple Data)은 하나의 Instruction으로 다수의 Data를 처리하는 아키텍처입니다. ☺️

 

(SIMD는 벡터 및 행렬 연산을 통해 비디오 인코딩이나 그래픽 렌더링 등에서 성능상 이점을 가져갈 수 있어요.)

 

 

 

 

여기서 Data란 64비트(MMX), 128비트(SSE), 256비트(AVX) 등의 크기를 가진 자료를 말합니다.

 

Arm은 Intel의 AVX, AVX512 등의 데이터를 지원하지 않지만

Neon이라 불리는 64-bit / 128-bit hybrid SIMD를 지원합니다.

 

 

 

 

 

암튼 왜 SIMD에 대하여 얘기했냐면

Metal은 <metal_simdgroup_matrix>에 정의된 

simdgroup_matrix<T,Cols,Rows> Matrix Type을 지원한다는 걸 말하려는 빌드업이었어요. ☺️

 

SIMD-group Matrix의 작업은 SIMD-group에 있는 스레드들에 의해 협력적으로 실행되기 때문에

모든 작업은 SIMD-group 내에서 균일한 제어 흐름에서만 실행되어야 합니다.

그렇지 않으면 동작이 정의되지 않아요. 🥺

 

 

 

 

 

 

 

Metal은 두 가지 모든 OS에 대하여 Metal 2.3 부터 simdgroup_matrix<T,Cols,Rows> Matrix Type을 지원합니다.

1. simdgroup_half8x8

2. simdgroup_float8x8

T는 half 또는 float이고, Cols와 Rows는 각각 8입니다.

 

 

 

 

 

SIMD-group의 스레드에 대한 행렬 요소의 매핑은 지정되지 않습니다.

SIMD-group에 지원되는 기능은 나중에 알아보겠습니다.

---

References

https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf

https://developer.arm.com/documentation/102159/latest/

https://en.wikipedia.org/wiki/Single_instruction,_multiple_data

https://github.com/neurolabusc/AppleSiliconForNeuroimaging