오랜만에 글을 써 봅니다. 원래 쓰던 CUDA 글도 마무리지었어야 했는데, Ubuntu의 병맛으로 여름방학식하고 다음날에 쓰던 글이 날아갔습니다(...) 그 이후로는 수능 공부도 하고 뭐 입시하다보니 글을 쓰지 않았습니다. 뭐... 일단 CUDA 글은 차차 마무리하기로 하고, 입시 끝난 지금부터는 새로운 기기를 다뤄보려고 합니다.

아, 참 이쁘지 않나요? 뭐 기능에 대해서도 끌리는 면이 있었지만, 일단 이쁘다는 것이 매우 끌렸...


  오래전부터 Xeon Phi에 대해 관심이 있었습니다. TOP500에서 한동안 1위를 했던 텐허-2도 Xeon Phi 기반의 슈퍼컴퓨터이고, 무엇보다 2016년즈음에 Xeon Phi 31S1P 버전이 170달러, 10개?넘게 구매하면 130달러인가까지 내려가는 상황도 있었기 때문이죠. 이때 구매를 했어야 했는데... 여건이 안 되서(관심은 있었지만, 그만큼의 돈을 들이기도 좀 부담스러웠을 것이고, 당시 쓰던 제 장비에서는 돌리기 힘들었죠. 물론 샀으면 진짜 ㅈ됐을 것입니다. Passive Cooling 어쨌을꺼...) 구매를 하지 않았습니다. 그 이후로 간간히 머릿 속에 떠오를 때마다 중고나라, 2CPU 등을 뒤져보면서 매물을 찾아보려 했지만, 애초에 있는 것이 이상한 법...(근데 중고나라에 팔렸던 전적이 있던 게 신기...) 없었죠.

  그러다가 이번에도 또 Xeon Phi 사고 싶어서 뒤지다가 ebay에 쳐보니 아니, 200달러에 5110P를 파네요? 개꿀하면서 하나 구매해서 굴려먹으려다가, 또 지인이 알려줘서 Xeon Phi 71S1P 6개(!!) 경매에 참여하게 되었습니다. 대충 Xeon Phi외에 서버 조립에 들일 금액은 100만원 정도 예상하고, 경매에는 최대 500달러까지 걸어두는 등, 엄청난(?) 짓을 했었는데 325달러에 낙찰받았습니다. 개당 54달러 정도에 얻은 편이니(물론 중고라고 해도) 꽤 싸게 구한 편이었...다고 생각은 했지만 이제부터 시작이었습니다.

  일단 배송비...USPS로는 받을 수 없으니 DHL를 쓰니 배송비 150?달러가 붙고, 관세도 내면서 대충 52?만원에 한국에 들여왔습니다. 자, 집에 왔습니다. 대충 수능 한 달 전쯤에 받은 거 같은데, 그 Intel의 푸-른 광채는 꽤나 이쁘더군요. 뭐 몇 개는 흠집이 있었지만, 그정도야 중고면 넘어 가아죠~


여기서 문제는...Passive Cooling라는 것. 일단은 암튼 idle 상태로라도 인식이 될 것이라 생각하고, 컴퓨터에 깔려 있던 Ubuntu 파티션을 줄이고 CentOS(Intel에서 권장하는 환경이 CentOS이고, Ubuntu에서도 설치할 수 있게 몇몇 분들이 해둔 것이 있었는데, 전 잘 안 돼서 그냥 바로 CentOS를 깔았습니다)를 설치하여 MPSS(Manycore Platform Software Stack, 네이밍 수준이...)를 설치했습니다. 대충 PCIe 인식은 잘 되었고, MPSS 서비스를 실행해서 잘 돌아가는지 확인만하면 되는데 안 되더라고요...? 여기서부터 이제 중고인데 어디 망가졌나... 싶었습니다. 일단은 수능이 끝나기 전까지 접어두기로 했습니다. 뭐 중간에 공부하기 싫어서 간단히 찾아보면서, 라이저 케이블, Xeon Phi용 쿨러(니덱 3열 쿨러를 샀었는데, 마이닝 케이스를 사면서 쓸모없어 졌습니다 ㅋ 5만원 ㅂㅂ), 마이닝 케이스(GPU를 많이 걸어두고 사용하니까, 딱 Passive Cooling을 쓰는 Xeon Phi에 적합하더라고요.) 등을 구매해두고 조립만 안 했습니다.

  자, 수능이 끝났습니다. 바로 조립하기 시작했습니다 ㅋㅋㅋ 일단 있는 걸로 대충 조립해서 Xeon Phi가 잘 돌아가는지 확인했습니다. 

예상대로 조립이 안 되어서, 결국 마이닝 케이스에 달려있던 쿨러 1열을 바깥으로 꺼내고 Xeon Phi를 철사로 묶어서 매달아 버렸습니다. 이후에는 철사 말고 케이블타이로 묶으니 완벽하더군요. 다행히 쿨러 다 작동시키니 정상적으로 MPSS 서비스가 돌아가면서 인식이 잘 되었습니다. 대충 케이스 질량이 7kg, Xeon Phi가 2개에 2kg?, 기타 잡다한 거까지 하니 10kg은 나온 거 같네요. 저걸 들고 옮기고, 허리 굽혀서 작업하느라 근육통 나서, 허리랑 팔, 다리 아파서 다음날 죽는 줄 알았습니다.

  중고로 메인보드 & CPU 같이 구매하면서 딸려 온 CPU 쿨러의 경우 히트싱크?가 너무 커서 이 위에 Xeon Phi를 올릴 수 없었습니다. 그래서 쿨러텍에서 뭐 3만원에 파는 싼 거 샀는데, 염병... 이게 뭐 소켓은 맞는데 쿨러 규격이 안 맞아서 조립이 안 되더군요. 반품하고 6만원짜리 Phanteks꺼 쿨러를 샀습니다. 작고 잘 돌아가네요.(지금 CPU를 풀로드할 일이 없어서 정확히는 완벽하다고는 못 말하지만, 뭐 암튼 잘 돌아가겠죠 ㅋㅋㅋ 어차피 소음이 큰 서버라 쿨러가 풀로드라도 잘 돌아가기만 하면 되니까) 암튼 이후에 6개 다 연결했습니다. 그리고 이게 마지막 모습입니다(...)

이게... 라이저 케이블이 한 2cm 정도 짧아서 원래 사 둔 테스트용 다른 케이블이랑 같이 섞어서 연결했는데, 중간에 뭐 접촉불량이든지 그런 게 있었는지, 전원을 넣자마자 테스트용 라이저 케이블에 달려 있던 콘덴서가 터졌습니닼ㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋㅋ 아 개웃기네. (당시에는 Xeon Phi 안에 있는 콘덴서 터진 줄 알고 욕하면서 서버 껐음...ㅠㅠ) 죽은 줄 알고 막 Xeon Phi 영정사진도 만들었는데, 다행히 다시 테스트해보니 인식은 되는 거 같습니다. Xeon Phi OS가 부팅되는 것까진 확인이 안 됐지만... 일단 Aliexpress에서 좀 더 긴 놈 주문했으니 이거 오면 연결해보려고 합니다. 그래서 지금은 5개 연결되어 있습니다. 이후에 이제 전원이 부족하니까(파워 서플라이에서 제공하는 6GPU용 케이블을 연결하면 깔끔하게 다 끼울 수 있으나, Xeon Phi가 각각 TDP가 300W인지라... 6개하고 본체 전원까지 생각하면 2kW는 생각해야죠) 같은 파워 하나 더 주문해서 듀얼 서플라이를 세팅했습니다. 모 분의 도움을 통해 MOSFET을 연결하여 동시에 켜지도록 구현했고요. 아직 무거운 프로그램을 돌려보지 않아서 잘 모르겠지만, 일단 기본적으로는 완벽하게 잘 돌아가네요. sample 코드도 잘 돌아갔고요. 문제는, idle 상태에서도 다 합쳐서 5개 기준으로 730W 정도 먹는다는 것. IDC에 넣어 두거나, 전기료 좀 싸게 쓸 수 있는 방법 찾기 전까지는 막 돌리다간 어머니한테 등짝 나갈 수도 있겠네요. 현재까지의 완성모습입니다.

(옆에 철사는 접지용 철사입니다. 이게 어디서 전류가 새는지 0.5V 정도 흐르더라고요. 접지했습니다.)


  이제, 제가 여행가기 전까지(2019년 1월 4일 전까지) Xeon Phi 프로그래밍하면서 개발일지와 흑우짓(...)을 올려보려고 합니다. 자료가 거의 없어서 거의 Intel 공식 문서 기반으로 할 듯... 목표는 OpenFOAM 돌려보고(Xeon Phi KNC 지원하는 것으로 앎) 직접 과학시뮬레이션 돌려보는 것입니다. 입시 끝났는데 딱히 게임도 재미가 없어서 이런 거나 하렵니다 ㅋㅋ


의문1. 도대체 콘덴서는 왜 터졌는가... 일단 Xeon Phi에 큰 영향을 주진 않은 거 같아서 다행인 거 같지만...

의문2. Colorful CROSSFEED 1250W "80PLUS Platinum"은 어떻게 그런 가성비가 나오는 건가...

'Programming > Intel® Xeon Phi™' 카테고리의 다른 글

0. Xeon Phi 서버 삽질과 세팅  (1) 2018.12.07
  1. ㅂㅅㅂ 2018.12.09 23:54 신고

    뽀대는 장난아니네요. 좋은 무기를 드셨으니즐거운 연산하세요. 잘 보고갑니다.

  블로그에 재미난 글을 써보려고 고민을 좀 하다가 문득 제가 요즘 즐기는 게임인 PLAYERUNKNOWN'S BATTLEGROUND(이하 배그)에서 나오는 여러 가지 현상을 물리학적으로 나름 엄밀하게(기준은 최소 고등학교 물I, 물II) 분석해보면 재미있을 거 같아서 새로운 주제의 글을 쓰게 되었습니다. 물I, 물II에서 재미없게 배운 물리학 지식을 이런 곳에서 적용해보면 나름 좋을 거 같아서 말입니다. 그래서 오늘 쓰는 글은 배그에서 가장 처음으로 만날 수 있는 물리학적 현상인 '종단 속도'입니다.

- 배그에서 낙하 속도는 최대 234km/h가 나온다. -

 

  물II에서 포물선 운동을 배울 때 자유낙하라는 것을 배우긴 하지만, 이때는 공기 저항을 고려하지 않기 때문에 단순히 연직방향으로 중력가속도만 작용해서 등가속도 직선 운동을 합니다. 하지만 현실 세계에는 공기 저항이 작용하여 어느 이상의 속력을 내지 못하도록 만들죠. 만약 물II에서 배운 것처럼 공기 저항 없는 자유 낙하가 현실 세계에서 성립했다면, 빗방울 맞고도 골로 가겠죠. 물II에서 배우진 않았지만, 현실 세계에서 공기 저항이 존재한다는 것은 누구나 아는 사실입니다. 그런데 과연 배그에 나오는 종단 속도 234km/h는 정말 현실 세계에서도 똑같이 나올까요?


  먼저 이를 물리적으로 엄밀하게 분석하려면 자유 낙하하는 신체에는 어떤 힘이 작용하는지 알아봐야 합니다. 먼저, 중력이 작용할테고, 물I 유체역학 부분에서 배운 부력도 작용하고, 공기 저항으로 인해 발생하는 항력(drag force)이 있을 것입니다. 각 힘을 $F_g, F_b, F_d$ 라 하면 아래와 같이 표현할 수 있습니다.

 

이제 수식으로 정리해서 종단 속도를 유도해내면 됩니다. 하지만 고등학교 과정에서 항력은 배우지 않았기 때문에 좀 찾아봤습니다. 항력 공식은 아래와 같습니다.

$$F_d  = -\frac{1}{2}\rho v^2 A C_d \hat{\mathrm{v}} $$

설명을 하자면, $\rho$는 유체의 밀도, $v$는 유체에 대한 물체의 상대 속도, $A$는 운동 방향에 수직인 평면에 대한 정사영 면적, $C_d$는 항력 계수, $\hat{\mathrm{v}}$는 속도의 방향을 나타낸 단위 벡터입니다.

 

이제 이들 공식을 이용해서 운동 방정식을 세워보면 다음과 같습니다.

$$F_d + F_b + F_g = ma$$

$$-\frac{1}{2}\rho v^2 A C_d \hat{\mathrm{v}} - \rho g V + mg = ma$$

이때 $a = 0$ 일 때 힘의 평형을 이루므로 종단속도를 가지게 됩니다. 따라서 식을 정리해보면 아래와 같습니다.

$$v = \sqrt{\frac{2\left(mg - \rho g V  \right)}{\rho A C_d}}$$

 

[2] 자료를 참고해서 상수를 아래와 같이 결정해보고 구해보겠습니다.

$m$: 70kg

$g$: 9.8m/s2

$C_d$: 0.7

$\rho$: 1 kg/m3

$A$:  0.18 m2(머리 방향으로 낙하할 때)

$V$: 0.075 m3 ([3]에서 인체의 평균 밀도 참고)

 

$$v = \sqrt{\frac{2\left( 70 \times 9.8 - 1 \times 9.8 \times 0.075\right)}{1 \times 0.18 \times 0.7}}  = 약\ 104.29 m/s = 375.444 km/h$$

* 사실 부력은 되게 작아서 없다고 보고 계산해도 됩니다.

 

흠 생각보다 더 빠르게 나오네요. [2]에 따르면 배를 아래로 했을 경우에는 200 km/h, 머리를 향했을 때 240~290 km/h가 나오고 항력을 최소한으로 했을 때가 480 km/h 나온다네요. 인체의 몸에 맞게 $C_d$를 좀 더 수정하면 실제 값에 가까워질 거 같네요.

 

  다음으론 단순히 종단 속도만 구하긴 아쉬우니까 어떻게 가속이 되어 종단 속도까지 도달하는지 분석해봅시다. 위에 있는 운동 방정식에서 a = 0으로 두지 않고 미분 방정식을 풀면 됩니다. 1계 비선형 미분방정식인 거 같네요. 어떻게 풀질 몰라서 WolframAlpha를 참고하면서 풀어봤습니다.


일단 아래와 같이 상수를 좀 줄여서 깔끔하게 한 상태로 방정식을 정리하고 풀어봤습니다.

$$ a = -\frac{1}{2}\rho A C_d,\ b = - \rho g V + mg,\  c = m \\ a v^2 + b = c \frac{dv}{dt}\\ 1 = \frac{c}{av^2 + b} \frac{dv}{dt}\\ \int\, dt = \frac{c}{b} \int \frac{1}{\frac{a}{b} v^2 + 1}\, dv \\ \sqrt{\frac{a}{b}} v = \tan u,\ \sqrt{\frac{a}{b}}\ dv = \sec^2 u\ du \\ \begin{align} \frac{c}{b} \int \frac{1}{\frac{a}{b} v^2 + 1}\ dv &= \frac{c}{b} \times \sqrt{\frac{b}{a}} \int \frac{\sec^2 u}{\tan^2 u + 1} \, du \\ &= \frac{c}{b} \times \sqrt{\frac{b}{a}} \int\, du \\ &= \frac{c}{b} \times \sqrt{\frac{b}{a}} \tan^{-1}\left(\sqrt{\frac{a}{b}} v\right) + C \end{align} \\\\  t = \frac{c}{b} \times \sqrt{\frac{b}{a}} \tan^{-1}\left(\sqrt{\frac{a}{b}} v\right) + C \\ \therefore v = \sqrt{\frac{b}{a}} \tan \left(\frac{\sqrt{ab}}{c}\left(t + C\right)\right)$$


이때, $t = 0$일 때 $v = 0$ 이므로, 속도는 아래와 같습니다.

$$v = \sqrt{\frac{b}{a}} \tan \left(\frac{\sqrt{ab}}{c}t \right)$$

(어째 tan 함수인게 불안한데...)


이제 위의 상수값을 대입하여 a, b, c의 값을 계산하고 정리하여 속도 방정식을 완성해보겠습니다.

$$a = -\frac{1}{2}\rho A C_d = -\frac{1}{2} \times 1 \times 0.18 \times 0.7 = -0.063 \\ b = - \rho g V + mg = - 1 \times 9.8 \times 0.075 + 70 \times 9.8 = 685.265 \\ c = m = 70 \\ \therefore v = -104.29 i \tan \left(0.09i\ t \right)$$

* 104.29i 가 아니라 -104.29i인 이유는 $a < 0, b > 0$에서 $\sqrt{\frac{b}{a}} = \sqrt{\frac{b}{-a}} \times \frac{1}{i} = - i \sqrt{\frac{b}{-a}}$ 이기 때문입니다. 이거 덕분에 함수가 음수 나와서 20분 동안 고민했네요;;; 역시 수1 같은 기초를 열심히 해둬야...


아무튼 이렇게 상수까지 다 때려박아보니 속도가 한없이 증가할 거 같이 생긴 tan 함수에 고2 이후론 거의 보지도 못한 i가 함수 안쪽과 바깥쪽에 등장하고 있네요. 뭔가 싶어서 WolframAlpha에 때려 넣어보니 tanh 함수로 바꿔줍니다. 따라서 아래와 같이 됩니다.

$$v = 104.29 \tanh \left(0.09\ t \right)$$


그려보면 아래와 같습니다. (x축이 시간(s)축, y축이 속도(m/s)축)


이렇게 해서 종단 속도도 구해보고, 시간에 따라 실제 낙하 속도가 어떻게 변하는지 확인해봤습니다. 다음 번에는 뭘 해볼지 고민 해봐야 겠습니다.

(정작 종단 속도 구하는 것보다 미방 푸는데 시간을 훨씬 많이 쓴 듯)


참고 자료

[1] https://ko.wikipedia.org/wiki/%ED%95%AD%EB%A0%A5

[2] https://en.wikipedia.org/wiki/Speed_skydiving

[3] https://en.wikipedia.org/wiki/Orders_of_magnitude_(density)

  1. Minjae Isaac Kwon 2018.04.01 14:00 신고

    F_r 이 v^2 = (dx/dt)^2 에 Dependent 하도록 풀면 (2계도) 비선형 미방이라 ♩♩♬맞습니다.
    그냥 근사 쳐서 v=dx/dt 에 맞게 풀면 일반해도 있고 리니어텀이라서 공간차원 추가해도 먹고 살기 편합니다.
    (2계도 2차 비선형 미방은 2차원으로 가는순간 수치해석밖에 노답)

  지난 번 글에서는 간단하게 왜 fractal 그리는 작업에 GPU가 필요한지에 대해 간단히 이야기를 해봤습니다. 이번에는 그 fractal을 그리는데 사용되는 CUDA에 대해 설명을 해보도록 하겠습니다.


  CUDA(Compute Unified Device Architecture)는 C언어(정확히는 CUDA C)를 사용해서 GPU에서 알고리즘을 작성할 수 있는 GPGPU 기술입니다. 정확히는 오직 Nvidia GPU에서만 지원하는 기술이죠. AMD GPU나 기타 다른 GPU에서도 비슷한 역할을 하는 OpenCL(Nvidia GPU에서도 돌아갑니다)이라는 것도 있지만, 제가 알기로는 CUDA가 더 좋다고 알고 있습니다. 요즘엔 딥 러닝 때문에 기본적으로 연산을 위해 GPU를 쓰는 경우가 있는데, 대체적으로 딥 러닝 프레임워크가 CUDA를 지원해서 딥 러닝 할 때는 어쩔 수 없이 Nvidia GPU를 고를 수 밖에 없는 일이 발생하기도 했습니다.


  원래 GPU는 CPU와 다르기 때문에 기본적으로 CUDA를 이용해 병렬 프로그래밍을 짤 때에도 기존 프로그래밍과는 조금 다른 방식으로 프로그래밍을 해야 합니다. 무엇보다 RAM을 동시에 같이 쓰지 않기 때문에 CUDA에서는 컴퓨터(Host)와 GPU(Device) 사이의 메모리 복사를 통해 작업물과 결과물을 주고 받습니다. 그리고 Kernel이라는 GPU에서의 function을 짜고 이를 GPU에서 돌려서 연산을 수행합니다. 즉, 기본적인 알고리즘은 아래와 같이 이뤄지게 됩니다.


1. CPU에서 GPU로 데이터 전송(Host to Device 메모리 복사)

2. GPU에서 kernel을 통해 데이터 연산을 수행(kernel 연산)

3. 연산 결과물을 GPU에서 CPU로 데이터 전송(Device to Host 메모리 복사)


이 과정을 원래는 CUDA의 기본 지원 언어인 C/C++에서 짤 때 일일히 구현해야 하는데, 제가 fractal 그리는 프로그램을 짤 때는 PyCUDA를 사용했기 때문에 저 과정을 간단하게 작성할 수 있었습니다. 제 프로그램에서 CUDA 부분만 가져와보면 아래와 같습니다.


kernel 부분

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
__global__ void calcMandelbrot(bool *isJulia, double *c, double *matrix, int *result) {
    int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_y = blockIdx.y * blockDim.y + threadIdx.y;
    double a = matrix[0+ (float)idx_x * matrix[2];
    double b = matrix[1+ (float)idx_y * matrix[2];
 
    double an = a, bn = b;
    double aan = an * an, bbn = bn * bn;
 
    for (int i = 0; i < 400; i+=5) {
        // iteration 1
        if (isJulia[0]) { // Julia Set
            bn = 2 * an * bn + c[1];
            an = aan - bbn + c[0];
            aan = an * an;
            bbn = bn * bn;
        } else { // Mandelbrot Set
            bn = 2 * an * bn + b;
            an = aan - bbn + a;
            aan = an * an;
            bbn = bn * bn;
        }
 
        if (an * an + bn * bn > 4.0f) {
            result[idx_x * gridDim.y * blockDim.y + idx_y] = i;
            break;
        }
 
        ... 생략 ...
 
    }
}
cs


main 부분

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
def calcMandelbrot(type, start, size, unit, c=0 + 0j):
    """ Calculating and drawing mandelbrot set and julia set
    :param type: fractal type
    :param start: center point in plane
    :param size: image size(width, height)
    :param unit: gap among pixels
    :param c: julia constant
    :return: image 2D array (format: RGBA)
    """
 
    s = time()
 
    # change size to multiple of 16(bigger or equal to prior size). Because of threads per block is 16
    size = ((size + 255/ 256).astype(np.uint32) * 256
 
    result = np.empty((size[0], size[1]), np.int32)
 
    func = cu.get_function("calcMandelbrot")
    func(cuda.In(np.array([type == Fractal.JULIA], np.bool)),
         cuda.In(np.array([c.real, c.imag], np.float64)),
         cuda.In(np.array([start[0- size[0/ 2 * unit,
                           start[1- size[1/ 2 * unit, unit], np.float64)),
         cuda.Out(result),
         block=(16161), grid=(int(size[0/ 16), int(size[1/ 16)))
 
    # Because in image symmetric transformation occurs between x axis and y axis
    result = np.transpose(result)
 
    return array2imgarray(result, cc.m_cyclic_wrwbw_40_90_c42_s25)
cs


main부분을 보시면 16번째 줄에서 결과물을 복사해서 담을 CPU 메모리를 할당하고 있고, 18번째 줄에서 function을 가져와서 19번째 줄에서 호출을 합니다. 원래 C/C++ 같았으면 19번째 줄에서 파라미터 부분(cuda.In, cuda.Out)을 싸그리 할당하고 복사하는 복잡한 과정을 거쳐야 호출할 수 있는데, PyCUDA에서는 간단하게 GPU에 입력만 될 파라미터는 cuda.In으로, GPU에서 출력만 될 파라미터는 cuda.Out으로 처리만 하면 나머지 할당하고 복사하는 과정을 알아서 해줍니다. 만약에 GPU에 입력된 후 다시 출력될 파라미터 같은 경우는 cuda.InOut으로 해주면 간단하게 처리됩니다.


그런데 잘 보시면 func 호출 시 단순히 파라미터만 있는 것이 아니라 block이나 grid와 같은 것이 보입니다. 이것은 GPU에 얼마 만큼의 병렬 처리 단위를 생성해서 계산을 할 것인지를 정하는 것입니다. 제가 겪었던 첫번째 시행착오가 바로 이 부분이었죠.


CUDA에서 병렬 처리 단위는 block과 thread로 나눠집니다. 기본적으로 여러 개의 block이 있고, 이 block 안에 thread가 있는 것입니다. 그리고 각 thread는 연산을 할 때 자신의 위치를 지표로 삼아 데이터를 불러와서 연산을 합니다. 예를 들어 fractal 이미지 같은 경우 픽셀의 (x, y)값마다 thread가 있을테니 각 thread는 자신의 (x, y)에 대한 값을 연산합니다. CUDA에서는 이 (x, y)와 같은 thread index가 3차원까지 되어 있으며 각 차원 별로, 그리고 각 단위(block, thread) 별로 제한값이 있습니다. 이것은 GPU마다 다른데 특히 block 당 thread의 최대 개수는 최근 device 경우에는 1024개로 제한이 되어 있습니다. 이 제한 값은 Linux에서 CUDA 설치 시 제공되는 Sample 중 deviceQuery라는 것을 컴파일해서 실행해보면 알 수 있습니다. 제 경우(GTX 970)엔 아래와 같이 확인할 수 있습니다.


GTX 970 Device Query


이렇게 각 block과 thread 개수를 정해서 GPU에서 kernel을 실행하면 개수에 맞게 thread가 생성이 되어 kernel 함수가 연산을 수행합니다. 이때 자기 thread가 어디 위치에 있는 것인지 알아야 하기 때문에 CUDA에서는 기본적으로 gridDim, blockIdx, blockDim, threadIdx라는 상수를 제공합니다. 제가 이걸 프로그래밍하면서 당연하게 'blockIdx는 block의 위치값, blockDim은 block의 차원일테니, blockDim으로 block의 총 개수를 보면 되겠다!'라는 생각을 했었는데, 이게 아니더라고요. block의 총 개수: gridDim, thread의 총 개수: blockDim 입니다. blockIdx는 block의 index, threadIdx는 thread의 index이고요. 이 내용을 간단하게 그림으로 나타내면 아래와 같습니다. 3D는 그리기 귀찮아서, 2D로 그려봤습니다.


  암튼 이렇게 CUDA의 기본 구조를 파악하고 나서 간단하게 fractal 그리기 프로그램을 만들었습니다. 초기 버전은 단순하게 block 당 thread를 1개로만 해서 만들었습니다. 이때는 뭐 거의 800x800 그리는데 0.3초 정도로 매우 느렸습니다. 그래서 어떻게 해야 하나 싶어서 CUDA sample을 뒤져서 Mandelbrot 예제(옛날에 CUDA sample 뒤지다가 발견했었거든요 ㅇㅇ)를 찾아내서 코드를 비교해봤습니다. 전체적인 구조는 비슷하더라고요. 차이점을 찾아보니 일단 저는 단일 thread만 사용했고, 이상하게 CUDA sample은 loop를 일부 풀어놨습니다(다음 글에서 소개할 기술입니다). 아무래도 thread가 하나인 게 문제인 거 같아서 thread를 여러 개로 늘려봤더니, 속도가 100배 정도(...) 빨라졌습니다. 위에서 말했다시피 thread는 개수가 겨우 1024개밖에 안 되기 때문에 몇 개일 때 최적의 값을 가질지 한 번 검색을 해봤더니, 알고리즘마다 case by case인 것 같더라고요. 제 알고리즘에서 이를 계산해보니 아래와 같은 그래프가 나왔습니다. 네, 정말, thread 한 개 쓰는 건 무식한 짓이라는 것이 너무 자명하게 보이죠. (지금 다시 그래프를 그려보니 100배까진 차이가 안 나네요, 처음 알고리즘을 도대체 어떻게 짰었길래...)

사실상 42 thread 정도 돌리면 이후는 거의 엇비슷한 속도가 나옵니다. 그 이후는 돌릴 때마다 누가 빠른지 조금씩 다르더라고요. 그래서 저는 fractal 그리는 프로그램은 162 thread로 작성했습니다.


이렇게 간단히 fractal 그리는 프로그램을 만들면서 기본적인 최적화를 해봤습니다. 그리고 CUDA sample에 나와있는 의미심장한 loop를 풀어 둔 것을 검색하면서 또 다른 것을 배웠죠. 다음 글에서 다뤄볼 예정입니다.


+ Recent posts

티스토리 툴바