main.cpp에서 CUDA 함수 호출해서 사용하기.

cmake를 이용해서 빌드할 때 main.cpp에서 cuda를 이용한 행렬곱 함수를 호출해 사용하는 예제이다.

결론부터 말하면 main.cpp는 matmul.cuh를 include하고 실제 쿠다를 사용하는 함수들은 모두 matmul.cu에 선언, 정의되어있다.
matmul.cuh에서 __global__ 등의 키워드를 가진 함수를 선언하거나 "cuda_runtime.h" 등을 인클루드하면 컴파일 에러가 발생한다.

정확한 동작은 알지 못하지만 cpp파일은 c++용 컴파일러, cu 파일은 nvcc를 이용해 컴파일하기 때문인 것 같다.

CMakeLists.txt

CMAKE_MINIMUM_REQUIRED(VERSION 3.26)
project(CPP_WITH_CU LANGUAGES CXX CUDA)

add_executable(main 
    src/main.cpp
    src/matmul.cu
)

target_include_directories(main PUBLIC
    include
)

include/matmul.cuh

#ifndef _MATMUL_CUH_
#define _MATMUL_CUH_

void matMulWrapper(int* A, int* B, int* C, int m, int n, int k);

#endif // _MATMUL_CUH_

src/main.cpp

#include "matmul.cuh"

#include <chrono>
#include <cmath>

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

int main(int argc, char **argv)
{
    int m, n, k;
    if (argc < 3)
    {
        m = 500;
        n = 600;
        k = 700;
    }
    else
    {
        m = atoi(argv[1]);
        n = atoi(argv[2]);
        k = atoi(argv[3]);
    }

    printf("matrix size A(%d, %d), B(%d, %d), C(%d, %d)\n", m, k, k, n, m, n);

    int *A = new int[m * k];
    int *B = new int[k * n];
    int *C = new int[m * n];
    int *C2 = new int[m * n];
    for (int i = 0; i < m * k; ++i)
    {
        A[i] = rand() % 100;
    }

    for (int i = 0; i < k * n; ++i)
    {
        B[i] = rand() % 100;
    }

    std::chrono::system_clock::time_point start = std::chrono::system_clock::now();
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < n; ++j)
        {
            int sum = 0;
            for (int d = 0; d < k; ++d)
            {
                sum += A[i * k + d] * B[n * d + j];
            }
            C2[n * i + j] = sum;
        }
    }
    std::chrono::duration<double> timeCpuMatmul = std::chrono::system_clock::now() - start;

    matMulWrapper(A, B, C, m, n, k);

    bool matrixCompare = true;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < n; ++j)
        {
            if (C[n * i + j] != C2[n * i + j])
            {
                printf("wrong value at (%d, %d) C1 = %d C2 = %d\n", i, j, C[n * i + j], C2[n * i + j]);
                matrixCompare = false;
            }
        }
    }

    printf("cpu matmul elapsed : %lf(ms)\n", timeCpuMatmul * 1000);

    if (matrixCompare == true)
    {
        printf("matmul results using CPU and GPU are matched\n");
    }

    delete[] A, B, C, C2;
    return 0;
}

src/matmul.cu

#include "matmul.cuh"

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <chrono>
#include <cmath>

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

__global__ void matMul(int* dA, int* dB, int* dC, int m, int n, int k)
{
	int row = blockDim.x * blockIdx.x + threadIdx.x;
	int col = blockDim.y * blockIdx.y + threadIdx.y;
	int index = row * n + col;
	if (row >= m || col >= n)
	{
		return;
	}
	int sum = 0;
	for (int i = 0; i < k; ++i)
	{
		sum += dA[row * k + i] * dB[n * i + col];
	}
	dC[index] = sum;
}

void matMulWrapper(int* A, int* B, int* C, int m, int n, int k)
{
    std::chrono::system_clock::time_point start = std::chrono::system_clock::now();
    int* dA, * dB, * dC;
	cudaMalloc(&dA, m * k * sizeof(int));
	cudaMemset(dA, 0, m * k * sizeof(int));
	cudaMalloc(&dB, k * n * sizeof(int));
	cudaMemset(dB, 0, k * n * sizeof(int));
	cudaMalloc(&dC, m * n * sizeof(int));
	cudaMemset(dC, 0, m * n * sizeof(int));
    std::chrono::duration<double> timeCudaMalloc = std::chrono::system_clock::now() - start;

    start = std::chrono::system_clock::now();
	cudaMemcpy(dA, A, m * k * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(dB, B, k * n * sizeof(int), cudaMemcpyHostToDevice);
	std::chrono::duration<double> timeCudaMemcpyToDevice = std::chrono::system_clock::now() - start;

	int blockSize = 32;
	dim3 gridDim(ceil(static_cast<float>(m) / blockSize), ceil(static_cast<float>(n) / blockSize));
	dim3 blockDim(blockSize, blockSize);
	printf("Grid(%d, %d), Block(%d, %d)\n", gridDim.x, gridDim.y, blockDim.x, blockDim.y);

    start = std::chrono::system_clock::now();
	matMul <<<gridDim, blockDim >>> (dA, dB, dC, m, n, k);
	cudaDeviceSynchronize();
    std::chrono::duration<double> timeCudaMatmul = std::chrono::system_clock::now() - start;

    start = std::chrono::system_clock::now();
	cudaMemcpy(C, dC, m * n * sizeof(int), cudaMemcpyDeviceToHost);
    std::chrono::duration<double> timeCudaMemcpyToHost = std::chrono::system_clock::now() - start;

    cudaFree(dA);
	cudaFree(dB);
	cudaFree(dC);

    printf("cuda malloc : %lf(ms), cuda memcpy to device : %lf(ms), cuda matmul : %lf(ms), cuda memcpy to host : %lf(ms)\n", timeCudaMalloc*1000, timeCudaMemcpyToDevice*1000, timeCudaMatmul*1000, timeCudaMemcpyToHost*1000);   
}

결과

큰 사이즈의 행렬을 곱하는 데 CUDA를 사용하는 것이 CPU를 이용한 것보다 몇배 더 빠른 결과를 얻었다.
다만 cudaMalloc에 대부분의 시간이 소요되는 것으로 보아 행렬곱을 여러번 하게된다면 미리 넉넉히 cudaMalloc을 해두고 그때그때 값을 memcpy해서 사용한다면 계산시간을 줄일 수 있을 것 같다.