[Metal] Performing Calculations on a GPU 정리

Basic Tasks and Concepts 중,  Performing Calculations on a GPU 정리

 

< 목차 >

 

1. 메탈 파일 만들고 MSL로 GPU 함수 작성하기

2. MTLDevice 준비하기
3. MTLCommandQueue 준비하기
4. 메탈 함수 불러와서 pipeline으로 컨버팅하기

5. MTLBuffer 만들고 버퍼에 데이터 load해주기

6. MTLCommand Buffer 만들기

7. MTLCommand Encoder 만들기

8. Command Buffer를 commit해서 commands 실행시켜주기

 

[1] MSL로 GPU 함수 작성하기

 

두 array의 elements들을 더한 값을 
result라는 array에 추가하는 코드를 살펴보자

 

아래 코드는 CPU에서 연산이 수행되고,  C로 쓰여졌다. 

Array addition, written in C

 

이 연산이 GPU에서 수행되게 하고 싶으면, 너는 이 함수를 Metal Shading Language (MSL)로 다시 작성해야한다. 

MSL은 GPU 프로그래밍을 위해 설계된 C ++의 변형이다. 

 

Metal에서는 GPU에서 실행되는 코드를 shader 라고 한다.

(이유는 역사적으로 3D 그래픽에서 색상을 계산하는 데 처음 사용 되었기 때문이다..?!??!?)

 

아래 코드는 위의 코드와 동일한 계산을 수행하는 MSL의 shader를 보여준다.
이 함수는 add.metal 파일에 선언되어있다. 

 

Array addition, written in MSL

이제 위의 함수를 자세히 살펴보자--!!

 

[1] kernel 키워드

 

위 함수는 kernel 키워드를 써서 두가지를 나타내고 있다.

 

1) public GPU function이다. 

public function은 다른 shader function들에 의해 불릴 수 없는 특성을 가지고 있다. 

(A public GPU function. Public functions are the only functions that your app can see. Public functions also can’t be called by other shader functions.)

2) grid of threads를 사용하여 병렬 연산을 하는 compute function 또는 compute kernel 이다. 

 

[2] device 키워드

 

세 개의 arguments에 device 키워드를 선언하고 있다. 

포인터들이 device address space에 있다는 것을 나타낸다. 

 

[3] for-loop 제거, thread_position_in_grid 키워드

 

이 함수는 compute grid에 있는 multiple threads에 의해 불릴 것이기 때문에 for-loop을 제거했다. 

배열의 각 element는 다른 쓰레드에 의해 계산된다. 

 

for-loop을 제거하기 위해

이 함수는 새로운 index argument를 만들었고, MSL keyword인 thread_position_in_grid를 선언했다. 

이 키워드는 메탈이 각각의 쓰레드에서 unique index를 계산하고 그 인덱스를 argument에 넘겨줘야한다 를 의미한다. 

 

 

 

===> 조금 어려운데, 그냥 add.metal 파일을 만들고 위의 함수를 복붙해주셔도 됩니당!!! 

 

 

 

[2] MTLDevice 준비하기

 

앱에서 MTLDevice 객체는 GPU에 대한 얇은 추상화이다. (thin abstraction)

이를 사용하여 GPU와 통신한다. 

 

MTLCreateSystemDefaultDevice() 로 default device를 얻을 수 있다.

(참고로 맥은 mulitple GPUs를 가지는 데, 메탈은 이 중 하나를 골라서 리턴해준다)

우리는 곧 device로 pipeline, commandQueue, buffer 같은 것을 만들 것이다. 

device.makeComputePipelineState
device.makeCommandQueue
device.makeBuffer

 

 

[3] MTLCommandQueue 준비하기

 

GPU로 작업을 보내려면 CommandQueue(명령 대기열)가 필요하다.

Metal은 CommandQueue를 사용하여 명령을 예약하기 때문이다. 

 

 

==>  여기까지하고 MetalAdder를 만들어보겠습니다..!! 

 

 

 

 

 

[4] 메탈 함수 불러와서 pipeline으로 컨버팅하기

 

Xcode는 모든 metal 파일들을 빌드해서 a default Metal library를 만들고 그것을 앱에 embed 시킵니다.

그래서 defaultLibary를 가져오고 거기서 메탈파일에 작성했던 함수를 가져올 수 있습니다..! 

 

하지만 이 함수 객체는 MSL function의 proxy라서 실행가능하지 않은 코드인데요,,,!

pipeline을 만들어서 이 함수를 실행가능한 코드(executable code)로 컨버팅 할 수 있습니다.

 

 

덧셈이라는 컴퓨팅 기능을 사용하기 때문에 MTLComputePipelineState를 만들었어요

 

 

 

 

 

 

 

여기까지 MetalAdder가 구현되었습니다-!

 

[5] MTLBuffer 만들고 버퍼에 데이터 load해주기

 

데이터를 들고 있을 버퍼를 만들어줍니다.

add_arrays 함수는 세 개의 배열이 필요해서 세 개 만들어주었습니다. 

 

 

 

 

 

 

참고로 copy 메소드 안쓰고 버퍼 만들때부터 데이터 넣어줄 수 도 있어요..! 

 

 

 

 

 

 

그리고 버퍼를 만들때, storage mode를 설정해줄 수 있는데, CPU 또는 GPU가 access 할 수 있는지?! 를 정해주는 것이라고 합니다.

샘플에서 사용한 storageModeShared는 CPU, GPU 모두 접근가능한 설정이라고 합니다. 

[6] MTLCommand Buffer 만들기

 

commands를 들고 있을 command buffer를 만듭니다.

 

 

[7] MTLCommand Encoder 만들기

 

이제 command buffer에 원하는 command를 써줘야합니다--!! 

 

command buffer에  command를 써주려면 command encoder를 사용해야합니다

샘플에서는 compute command encoder를 쓰는데, 이 인코더는 compute pass를 인코딩 합니다.

compute pass는 compute pipeline에서 실행되는 commands들의 목록을 hold 합니다.

 

 

 

 

 

commandEncoder를 만들고 

pipeline과 pipeline으로 전달할 arguments를 세팅해줍니다.

 

 

이렇게 해준 것입니다...!! 참고로 PSO는 pipeline state object를 의미합니다.

 

 

 

[8] Command Buffer를 commit해서 commands 실행시켜주기

 

commands를 실행시켜주기 위해서는 Command Buffer를 commit 해야합니다.

 

 

 

 

 

 

 

[9] 마지막으로 연산이 잘 수행되었는지 확인해주는 프린트문 추가하기

 

연산이 끝날때 까지 기다린 후, result 배열을 출력해서 덧셈 연산이 잘 되었는지 확인합니다-!

 

 

 

 

 

 

[최종코드]

 

 

 

 

 

 

 

 

 

 

 

(시뮬레이터 설정 따로 안해서 디바이스에서 돌려주세요-!)