[CUDA] 기본 정리
CUDA2022. 11. 1. 15:02
GitHub - NVIDIA/cuda-samples: Samples for CUDA Developers which demonstrates features in CUDA Toolkit
Samples for CUDA Developers which demonstrates features in CUDA Toolkit - GitHub - NVIDIA/cuda-samples: Samples for CUDA Developers which demonstrates features in CUDA Toolkit
github.com
1. include 및 개요
#include "cuda_runtime.h"
#include "device_launch_parameters.h"- 문법은 C/C++과 거의 동일하지만, kernel 내에서 Class나 math.h내 함수 사용 같은 건 불가능했다. 대신 쿠다에서 지원하는 math 함수 목록이 있다.
- 과정: Device set → memory allocation 및 copy → add kernel(GPU로 돌릴 함수 등록) → 완료될 때까지 대기 → 연산 결과를 CPU로 memcpy → 메모리 할당 해제
- cuda library 함수 이름 형식:
cuda$(FunctionName) - return value:
cudaError_t(typedef int)cudaSuccess와 return된 status 값을 비교하여 성공/실패 확인
2. code
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#define SIZE 1023
#define BLOCK_CNT 1
#define THREAD_CNT SIZE
bool isCudaError(cudaError_t status, uint8_t* cudaArr, uint8_t* cudaDest)
{
auto ret = status != cudaSuccess;
if(ret)
{
cudaFree(cudaArr);
cudaFree(cudaDest);
}
return ret;
}
__global__ void my_func(const uint8_t* src, uint8_t* dest)
{
const unsigned int index = blockIdx.x * THREAD_CNT + threadIdx.x;
// gpu로 돌릴 코드 작성
dest[index] = src[0] + src[4] + src[9] + src[index] * 2;
// ...
}
int cudaTest()
{
uint8_t arr[SIZE] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, uint8_t dest[SIZE] = {};
uint8_t *dev_arr = 0, *dev_dest = 0;
cudaError_t status;
// device setting
status = cudaSetDevice(0);
if(isCudaError(status)) return status;
// kernel 내에서 사용할 배열 malloc
status = cudaMalloc((void**)&dev_arr, SIZE * sizeof(uint8_t));
if (isCudaError(status)) return status;
status = cudaMalloc((void**)&dev_dest, SIZE * sizeof(uint8_t));
if (isCudaError(status)) return status;
// cpu에서 gpu로 전달해줄 값이 있을 경우 memcpy
status = cudaMemcpy(dev_arr, arr, SIZE * sizeof(uint8_t), cudaMemcpyHostToDevice);
if (isCudaError(status)) return status;
// gpu로 돌릴 함수 등록
my_func<<<BLOCK_CNT, THREAD_CNT>>> (dev_arr, dev_dest); // 이 부분 문법이 C/C++과 다르다.
if (isCudaError(cudaGetLastError())) return status;
// device sync
// cuda는 원래 비동기식이고,
// 이 함수를 사용하면 동기가 보장된다고는 봤는데 더 공부가 필요할듯
status = cudaDeviceSynchronize();
if (isCudaError(status)) return status;
// gpu로 연산한 값을 cpu로 memcpy
status = cudaMemcpy(dest, dev_dest, SIZE * sizeof(uint8_t), cudaMemcpyDeviceToHost);
if (isCudaError(status)) return status;
// cudaMalloc으로 할당한 메모리 해제
cudaFree(dev_arr);
cudaFree(dev_dest);
return status;
}- 예전에 작성했던 코드를 참고해서 손코딩한 코드이므로 안 돌아갈 수도 있다. 대략적인 흐름 참고용. 실제로 작업했던 코드는 https://github.com/temphi20/cuda-test를 참고.
- gpu(kernel) 코드 동작
- n개의 block 안에서 m개의 thread로 돈다. 1블럭당 최대 thread 수는 1024개이며, 이 수를 넘을 시 에러를 반환한다.
- 변하지 않는 값을 전달할 경우 함수 parameter type를 const로 적어주는 편이 효율적이다.
- 아마 block 순서대로 프로그램이 굴러가는 것 같다. (block 수를 최소화하도록 설계) ← 함수 내 프린트 결과로 추측한 부분이라 사실 확인 필요
- thread는 병렬로 진행하므로, 연속된 작업일 경우 한 thread 내에서 해결하도록 한다.
- gpu로 돌아갈 함수 앞에는
__global__을 붙인다. 작성하는 함수는 한 thread 내에서 하는 일에 해당된다. - 순차적인 index 접근을 하는 게 아니기 때문에(병렬 처리) 배열 접근시 index 계산은 코드 작성자가 알아서 해줘야 한다. 테스트 시
blockIdx.x와threadIdx.x로 접근했다. - 전달하는 변수도 지나치게 많거나 크고 block당 thread 갯수도 많을 경우 메모리 부족인가...싶은 오류가 발생할 수 있는데, 어지간하면 잘 안 생기는 듯. 이 오류가 발생했을 때도 몇 억 개인가 몇 십억 개인가, 한꺼번에 돌렸었다.
3. CUDA를 이용한 라이브러리 만들기
- window 환경 테스트는 해보지 않았지만, gcc를 이용할 경우 기존 C/C++ makefile 작성법과 크게 다르지 않았다.
- 단, 컴파일러 옵션으로
-Xcompiler를 추가해줘야 한다. - https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/#options-for-passing-specific-phase-options 참고.
4. 함수 목록
- module: https://docs.nvidia.com/cuda/cuda-driver-api/modules.html#modules
- math: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#standard-functions
5. 겪었던 에러 모음
- [100] no cuda-capable device is detected
- CUDA 지원이 안 되는 경우, 혹은 nvidia driver를 인식하지 못하거나 driver에 문제가 생긴 경우.
- cuda package 및 driver 간 version이 달라서 생기는 문제일 수도 있다.
- [804] forward compatibility was attempted on non supported HW
- 버전 문제로 추측.
- nvidia driver가 업데이트 실패해서 알림이 뜨고 있었는데,
sudo apt-get install -f로 실패한 업데이트를 설치해준 후 재부팅하니까 해결되었다.
- [700] an illegal memory access was encountered
- GPU 연산 중 허용되지 않은 메모리에 접근했을 경우 발생.
- 할당한 배열의 범위 밖에 해당하는 index에 접근했을 때 해당 오류가 발생되었다.
- 코드 작성만 잘 하면 안 일어나는 오류이므로 꼼꼼한 작성 및 디버깅 필요.
댓글()
