Jetson에서 cudaMallocManaged()는 CPU와 GPU가 같은 포인터로 접근할 수 있는 CUDA 통합 메모리를 할당한다.

Jetson은 CPU와 GPU가 물리 메모리를 공유하지만 동기화가 사라지는 것은 아니다.

형태는 cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags) 이다.

할당한 포인터를 CPU 코드와 CUDA 커널에서 모두 사용하고 cudaFree()로 해제한다.

커널이 끝나기 전에 CPU가 결과를 읽으려면 반드시 스트림 또는 장치 동기화가 필요하다.

아래는 커널 정의와 오류 검사를 포함한 CUDA C++ 예제다. 일부 Jetson은 프리페치와 호스트 등록 지원 조건이 다르므로 데스크톱 예제의 최적화 API를 무조건 추가하지 않는다.

#include <cuda_runtime.h>
#include <cstdio>
#include <cstdlib>
#define CHECK(call) do { \
    cudaError_t status_=(call); \
    if (status_!=cudaSuccess) { \
        std::fprintf(stderr,"%s\n",cudaGetErrorString(status_)); \
        std::exit(1); \
    } \
} while (0)
__global__ void fill_kernel(int *values, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count)
        values[i] = i;
}
int main() {
    int *values = nullptr;
    CHECK(cudaMallocManaged(&values, 1024 * sizeof(*values)));
    fill_kernel <<< 4, 256>>>(values, 1024);
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
    std::printf("%d %d\n", values[0], values[1023]);
    CHECK(cudaFree(values));
}

통합 메모리는 코드를 단순하게 하지만 모든 상황에서 가장 빠른 방식은 아니다.

페이지 폴트와 CPU/GPU 접근 전환 비용이 생길 수 있어 접근 패턴을 생각해야 한다.

Jetson에서 빌드하기

nvcc -std=c++17 main.cu -o main
./main