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에 접근했을 때 해당 오류가 발생되었다.
  • 코드 작성만 잘 하면 안 일어나는 오류이므로 꼼꼼한 작성 및 디버깅 필요.

다운로드 링크는 아래 참조. 하라는대로 하면 깔리기 때문에 크게 어려운 건 없다.

CUDA 지원 그래픽카드 확인은 아래 링크에서. CUDA 11.7은 nvidia-driver-515 이상부터 호환이 된다고 한다.

겪었던 에러 및 시행착오, 팁:

• linux runfile (local) 방식을 선택했을 경우 설치 실패
  • err log 확인 결과: driver 설치 실패
  • 해결: gpu driver를 제외하고 설치, 드라이버는 apt-get으로 따로 설치해주었다.

sudo apt-get install nvidia-driver-$VERSION
sudo reboot

• driver까지 모두 설치했는데 linux가 driver를 인식하지 못함
  • nvidia driver 설치 후, reboot할 때 MOK 등록을 하지 않아서...
  • Perform MOK management 화면(파란 화면) → [Enroll MOK] → [View key 0] → [Esc] → [Continue] → [OK] → [Reboot] 과정을 거치니 문제 없이 인식됨.
  • 참고했던 링크 목록:

• 환경변수 추가
  • window: 설치시 자동으로 추가된다.
  • linux: /usr/local/cuda-$(VERSION)/bin에 위치해있다. 이 경로로 등록하면 끝.
• linux에서 nvidia driver가 제대로 인식되고 있는지 확인하는 간단한 방법
  • nvidia-smi 명령어 사용.
  • 아래처럼 뜨면 제대로 되고 있는 거다.