VSCode 에서 CUDA (nvcc) 개발 환경 설정하기 – 1

이전 포스팅에서 그래픽 카드 붙은 시스템까지는 구성했고 이어서 개발환경 구축을 위한 첫걸음. 사실 NVIDIA Container Runtime 만 구성하면 NVIDIA NGC 에서 CUDA 뿐만 아니라 여러 라이브러리, 프레임워크 버전별 이미지를 제공해주다보니 직접 버전별로 구성해놓을 필요가 없어서 시스템을 깔끔하게 관리할 수도 있고 집에 있는 개인 장비, Cloud Workspace 그리고 폐쇄망인 회사 업무 환경에 까지 일관된 개발환경을 구성할 수 있을 것 같다.

VSCode 에서 Docker 와 Dev Conatainers 확장만 이용해도 제법 그럴싸할 것 같은데 일단 NGC 에서 아무 개발 이미지를 하나 대충 받아다가

vim docker-compose.yml

version: '3'

services:
  nvcc:
    image: nvcr.io/nvidia/cuda:11.3.0-devel-ubuntu20.04
    command: nvcc matmul.cu
    working_dir: /src
    volumes:
      - ./src:/src

무려 8년 전, 학부시절 만들어놓은 matmul 코드를 돌려봤는데

vim src/matmul.cu

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda.h>

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int width) {
    int tx = blockIdx.x * blockDim.x + threadIdx.x;
    int ty = blockIdx.y * blockDim.y + threadIdx.y;
    float Pvalue = 0;
    for (int k = 0; k < width; ++k) {
        float Mdelement = Md[ty * width + k];
        float Ndelement = Nd[k * width + tx];
        Pvalue += Mdelement * Ndelement;
    }
    Pd[ty * width + tx] = Pvalue;
}

void MatrixMultiplication(float* M, float* N, float* P, int width) {
    int size = width * width * sizeof(float);
    float *Md, *Nd, *Pd;

    cudaMalloc((void**) &Md, size);
    cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
    cudaMalloc((void**) &Nd, size);
    cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
    cudaMalloc((void**) &Pd, size);

    dim3 dimBlock(16, 16);
    dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x,
                 (width + dimBlock.y - 1) / dimBlock.y);

    MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, width);
    cudaDeviceSynchronize();

    cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

    cudaFree(Md);
    cudaFree(Nd);
    cudaFree(Pd);
}

int main() {
    clock_t start = clock();
    printf("program start\n");

    int width = 1000;
    float *M = (float*)malloc(width * width * sizeof(float));
    float *N = (float*)malloc(width * width * sizeof(float));
    float *P = (float*)malloc(width * width * sizeof(float));

    for (int i = 0; i < width * width; i++) {
        M[i] = (i % 2) ? 0.0f : 1.0f;
        N[i] = (i % 2) ? 1.0f : 0.0f;
    }

    MatrixMultiplication(M, N, P, width);

    printf("program finished\n");
    double duration = (double)(clock() - start) / CLOCKS_PER_SEC;
    printf("%f sec\n", duration);

    free(M);
    free(N);
    free(P);

    return 0;
}

이 프로젝트랑 같이 있던 Makefile 에서, 당시 최신 버전이었을, cuda-7.0 쓰도록 명시해놨길래 분명 컴파일이 안 될 줄 알았는데 너무 기본 api 만 사용하는 프로젝트여서 그런지 신기하게도 잘 돌아간다.

# docker compose up
[+] Running 1/0
 ⠿ Container matmul-nvcc-1  Created                                                                           0.0s
...
matmul-nvcc-1 exited with code 0
# src/a.out 
program start
program finished
0.172547 sec

이어서 VSCode가 C/CUDA IDE 로 얼마나 쓸만할지 이것저것 확장들을 조합해 개발 환경을 구축해보고 토치 같은 프로젝트 뜯어볼 수 있는 환경까지 만들어봐야겠다. 더 나아간다면 rust 개발환경까지 구축해서 PyO3 공부해보기.