https://www.coursera.org/ 에서 제공하는 강의를 공부하는 중인데요. 참 재미있네요.
요즘에 보고 있는 건, Heterogeneous Parallel Programming, 한국어로 하면 병렬 프로그래밍이죠.
이해하기로는 그래픽카드 제조업체인 NVidia에서 제공하는 CUDA 언어를 이용해서 병렬 프로그래밍을 배우는 건데요. 그래픽카드 제조업체이다보니, 행렬 계산을 하는 프로세스에 대해 더 많이 다뤘고, 그래서 그 쪽에 더 적합한 라이브러리를 제공할 수 있었던 것 같습니다.
NVidia와 CUDA에 대해서는 http://www.nvidia.co.kr/object/cuda-kr.html를 참고하시면 될 것 같습니다.
강의는 비디오 강의와 퀴즈, 프로그래밍 실습으로 구성되어 있는데요, 이번에 블로그에 올릴 글은 제가 프로그래밍 실습 과정에서 했던 삽질을 설명해볼까 합니다.
첫번째 주의 MP0는 너무나 쉽습니다. 코딩되어 있는 걸 그냥 웹상에서 실행하면 끝입니다.
그런데 문제는 MP1이었죠.
영어가 워낙 딸려서, 뭘 어떻게 하라는 건지 알 수가 없더군요. (아니면 강의를 성실히 듣거나요. --;; 강의 대강 듣고 프로그래밍하면서 배우려고 하면 저처럼 고생합니다. --;;)
일단 하라는 건 두개의 벡터를 더하라는 거였는데요. c에서는 그냥 for문 쓰면 될 일인데, 그걸 어떻게 하라는 건지. --;;
그래서 웹도 훑어보고, 검색도 해보고 해서 겨우 어떻게 하라는 건지 이해했습니다. 과정은
(강의에 나온 바와 같이)
1. 자료를 호스트(그러니까 컴퓨터 메모리?)에 읽어온다.
2. GPU에서 쓸 메모리를 배정한다.
3. 읽어온 자료를 GPU에서 사용하는 메모리로 복사한다.
4. GPU에서 사용할 블럭(BLOCK)과 그리드(GRID)의 크기를 정한다.
5. kernel 함수를 이용해서 CUDA로 계산한다. (GPU가 사용되겠죠?)
6. GPU 메모리에 있는 계산된 값을 호스트(CPU에서 사용하는 메모리, 그러니까 컴퓨터 본체 메모리)로 복사해온다.
7. GPU에서 사용했던 GPU 메모리를 돌려준다. (자꾸 쓰고 안 돌려주면 메모리가 줄어들겠죠.)
MP1에서 해야될 프로그래밍의 코드 중, 위의 1번은 이미되어 있습니다. 친절하죠? 이거 없었으면 아예 좌절했을 지도 모릅니다.
위의 2번.
CUDA에서는 메모리를 배정할 때 cudaMalloc 함수를 사용합니다.
저는 처음에 조금 망설였던게, 위 1번에서 자료를 읽어올 때에는 "hostInput1 =" 이렇게 되어 있어서, "deviceInput1 =" 이렇게 해야되나? 고민했는데요. 그건 아니고, cudaMalloc() 안에 GPU 메모리 값과 크기만 지정해주는 걸로 됩니다. 자세한 건 cudaMalloc 함수를 검색해보시면 더 자세하게 알 수 있을 겁니다. 메모리 값을 쓰실때에는 '&'를 써주시면 됩니다. 검색해본 결과 (void**)를 사용하는 예도 있는데요. 꼭 써야되는지는 잘 모르겠습니다.
위의 3번.
메모리를 배정해줬으니(사무실 임대를 했으니) 자료를 카피해줘야죠 (이사해야죠).
자료카피는 cudaMemcpy함수를 씁니다. 검색해보시면 아시겠지만, cudaMemcpy( <복사할 곳>, <원 소스>, <크기>, <구체적인 방법> )으로 사용하면 됩니다.
위의 4번.
CUDA에서는 BLOCK과 GRID, THREAD 개념이 매우 중요한 것 같습니다. 문제는 저도 잘 모른다는 거죠. ㅋㅋ
잘 몰라서 위험하긴 하지만, 대강 이해하기로는 GRID > BLOCK > THREAD같구요. 아파트로 치면, GRID는 '단지', BLOCK은 '동', THREAD는 '호실'을 의미하는 것 같습니다. 각각의 '호실'에 자료를 넣어놓고 통으로 더하기 빼기를 하는 거겠죠.
어쨌든 CUDA에서 허용하는 블럭 사이즈 이내로만 설정해주면 되는 것 같습니다. 참고로 전 BLOCK 사이즈로 256을 줬습니다. 어떤 예에서는 1000이상을 주기도 하더군요.
그 다음엔 그리드 크기를 지정해줘야 합니다. 이 그리드 크기는 자료 전체 양과 BLOCK의 크기에 연동되어서 지정해주게 되는 것 같더군요. 간단히 얘기해서 자료양이 BLOCK보다 몇배가 되느냐, 정도로 생각됩니다. 그런데 주의해야할 점은 integer값을 줘야 한다는 겁니다. 그래서 ceil함수를 써주는 게 좋습니다.
위의 5번.
자 이제 준비가 다 되었습니다. GPU메모리에 자료도 넣어놨고, 크기 할당할 값도 마련되었습니다. 그럼 이제 계산만 하면 되죠.
이 부분이 병렬계산의 묘미가 아닌가 싶습니다. 만일 for문으로 계산하려면, i,j,k,...을 쓰고, 중간에 헷갈리면 무한루프 돌기도 하고. ㅋ
그런데 CUDA에서는 간단합니다. 각 블럭의 주소끼리만 더하기 해주면 GPU가 알아서 두 블럭을 더해버립니다. for를 쓸 일이 없죠.
이 때 중요한 게, GPU 메모리 주소를 잘 지정해줘야 되는데요. 강의에도 나오지만, "blockIdx.x * blockDim.x + threadIdx.x"로 id를 지정해주고, 더하기 해주면 바로 됩니다. 자세한 코드는 검색해보시면 됩니다.
위의 6번.
계산이 끝났으니, 결과물을 배정한 GPU 메모리 (deviceOutput)에는 결과물이 들어있을 겁니다. 이제 그걸 CPU 메모리로 복사해옵니다. 물론 복사는 cudaMemcpy를 쓰면 됩니다. 위에 GPU로 복사해왔던 것과 정반대로 하면 되죠.
위의 7번.
GPU 메모리를 다 썼으니 이제 돌려줘야죠. cudaFree()를 써서 되돌려줍니다.
예제를 하면서 헤맸던 부분은,
- 메모리 크기 설정
; 전체 자료양과 그 자료의 형식(int인지, double인지, float인지)이 기본적으로 가져야할 메모리 크기에 주의
- 블럭과 그리드 크기 설정
; 처음엔 개념이 잘 안 서는데, 그냥 허용치 이내 값만 지정해주면 될 듯. 그리드 크기 설정 시 int값이 제대로 되도록 조심
- kernel함수를 불러올때 <<< >>>를 함수뒤에, 즉, "vecAdd<<<>>>();"와 같이 써야함. 그것도 3개씩.
; < 는 3개. ㅋㅋㅋ
입니다.
댓글 없음:
댓글 쓰기