DGX Spark CUDA 산출물 검증: GPU 대상 아키텍처와 PTX

DGX Spark에서 CUDA 프로젝트가 정상적으로 빌드되지만 커널을 실행할 때 no kernel image is available for execution on the device 오류가 발생한다면, 환경 전체를 재설치하기보다 먼저 산출물에 어떤 GPU 아키텍처가 포함되어 있는지 확인해야 합니다. 더 알아차리기 어려운 경우도 있습니다. 프로그램은 계속 PTX JIT 컴파일에 의존해 실행되는데, 릴리스 기록에는 네이티브 디바이스 코드가 포함된 것으로 잘못 기록되어 있는 경우입니다. 이 글에서는 독립적인 소형 프로그램으로 대상 조회, 산출물 검사, 두 실행 경로 검증을 하나의 인수 검증 절차로 연결합니다.

1. 호스트 아키텍처와 GPU 대상 구분하기

DGX Spark는 NVIDIA GB10 Grace Blackwell을 사용하며, 호스트 아키텍처는 ARM64입니다. 운영체제는 Ubuntu 24.04 기반의 NVIDIA DGX OS이며 CUDA 스택이 사전 설치되어 있습니다. 결제 후 콘솔에서 접속 정보를 확인하고 SSH로 작업할 수 있습니다.

ARM64는 CPU 측 실행 파일의 아키텍처를 나타내고, CUDA의 연산 능력(Compute Capability)은 GPU 디바이스 코드의 대상을 나타냅니다. 둘은 서로 다른 호환성 조건입니다. file 출력에 AArch64가 표시되어도 호스트 측 형식이 올바르다는 뜻일 뿐, 실행 가능한 커널이 내부에 포함되어 있다는 증거는 아닙니다.

SASS는 대상 GPU용 기계어 코드이고, PTX는 드라이버가 JIT 컴파일할 수 있는 중간 표현입니다. 배포 패키지에 둘 다 포함할 수 있지만, 각각 따로 검사해야 합니다. 일반 실행이 성공했다는 이유만으로 두 경로가 모두 검증되었다고 판단해서는 안 됩니다.

아래 실험은 새로 만든 디렉터리 안에 소스 코드와 실행 파일만 생성하며, 드라이버나 시스템 툴체인, 업무용 프로그램은 변경하지 않습니다. 예제 결과는 호환성 인수 검증에만 사용하며, 성능 테스트 결과가 아닙니다.

2. 아키텍처 번호를 추측하지 말고 실제 연산 능력 조회하기

먼저 별도의 작업 디렉터리를 만들고 시스템 및 툴체인 정보를 저장합니다. nvidia-smi에 표시되는 CUDA 버전은 드라이버가 지원하는 범위를 나타내며, 현재 nvcc의 툴킷 버전과 같다는 뜻은 아닙니다. 따라서 두 출력 모두 보관해야 합니다.

mkdir -p "$HOME/cuda-artifact-check"
cd "$HOME/cuda-artifact-check"
uname -m > host-arch.txt
nvidia-smi > driver.txt
nvcc --version > toolkit.txt
nvcc --list-gpu-code > supported-targets.txt

다음 내용을 probe.cu로 저장합니다. 이 프로그램은 CUDA Runtime을 통해 디바이스 0의 연산 능력을 조회합니다. 실패하면 0이 아닌 종료 코드를 반환하며, 잘못 사용될 수 있는 대상 번호를 출력하지 않습니다.

#include <cuda_runtime.h>
#include <cstdio>

int main() {
    cudaDeviceProp p{};
    cudaError_t e = cudaGetDeviceProperties(&p, 0);
    if (e != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e));
        return 1;
    }
    std::printf("%d%d\n", p.major, p.minor);
    return 0;
}

nvcc probe.cu -o probe
CC=$(./probe) || exit 1
printf '%s\n' "$CC" > compute-capability.txt
grep -Fx "sm_${CC}" supported-targets.txt || exit 1

여기서는 의도적으로 GB10의 대상 번호를 하드코딩하지 않습니다. 조회한 sm_ 대상이 툴킷 목록에 없다면 먼저 툴킷의 지원 범위를 확인해야 합니다. 임의로 이전 번호로 바꿔 릴리스를 진행해서는 안 됩니다. 조회 프로그램이 실패하면 디바이스 접근이나 런타임 라이브러리 문제부터 해결하고, 컴파일 단계로 넘어가지 마세요.

3. SASS와 PTX를 명시적으로 생성하고 검사하기

다음 내용을 check.cu로 저장합니다. 스레드 하나만 실행해 정수를 기록한 뒤 호스트로 복사하여 결과를 확인하는 프로그램입니다. 커널 실행이 실패했는데도 성공으로 출력하는 일이 없도록 모든 CUDA 호출의 오류를 검사합니다.

#include <cuda_runtime.h>
#include <cstdio>

#define CUDA_OK(call) do { \
    cudaError_t e = (call); \
    if (e != cudaSuccess) { \
        std::fprintf(stderr, "%s\n", cudaGetErrorString(e)); \
        return 1; \
    } \
} while (0)

__global__ void write_value(int* p) {
    *p = 7;
}

int main() {
    int* device = nullptr;
    int host = 0;
    CUDA_OK(cudaMalloc(&device, sizeof(int)));
    write_value<<<1, 1>>>(device);
    CUDA_OK(cudaGetLastError());
    CUDA_OK(cudaDeviceSynchronize());
    CUDA_OK(cudaMemcpy(&host, device, sizeof(int),
                       cudaMemcpyDeviceToHost));
    CUDA_OK(cudaFree(device));
    if (host != 7) return 2;
    std::puts("PASS");
    return 0;
}

같은 작업 디렉터리에서 다음 명령을 실행해 네이티브 코드와 PTX를 모두 생성하도록 명시적으로 지정합니다.

CC=$(cat compute-capability.txt)
nvcc -O2 \
  -gencode "arch=compute_${CC},code=[sm_${CC},compute_${CC}]" \
  check.cu -o check

file check > binary-format.txt
cuobjdump --list-elf ./check > elf-list.txt
cuobjdump --dump-sass ./check > sass.txt
cuobjdump --dump-ptx ./check > ptx.txt

ELF 목록과 SASS 출력에 해당 대상 및 커널이 있는지, PTX 출력에 .version, .target, 커널 진입점이 포함되어 있는지 확인합니다. 파일이 비어 있거나, 대상이 일치하지 않거나, 예상한 커널이 없다면 검증을 통과시켜서는 안 됩니다. 명령의 종료 코드만 확인해서는 부족합니다.

중간 오브젝트가 아닌 최종 파일 검사하기

실제 프로젝트에는 정적 라이브러리 링크, 디바이스 링크, 패키징 등의 단계가 포함될 수 있습니다. 오브젝트 파일에 PTX가 있다고 해서 최종 납품 파일에도 남아 있다는 뜻은 아닙니다. 최종 실행 파일과 실제로 로드되는 공유 라이브러리를 대상으로 검사를 반복하고, 파일 해시를 기록해야 합니다.

4. 네이티브 실행과 JIT 컴파일을 각각 검증하기

먼저 PTX JIT 컴파일을 금지하여 현재 산출물이 네이티브 디바이스 코드로 실행되도록 합니다. 그런 다음 PTX 경로를 강제하고 디스크 컴파일 캐시를 비활성화합니다. 두 테스트 모두 새 프로세스에서 실행하여 프로세스 안에 이미 로드된 모듈이 판단에 영향을 주지 않도록 합니다.

env -u CUDA_FORCE_PTX_JIT \
  CUDA_DISABLE_PTX_JIT=1 \
  ./check > native.log 2>&1
native_rc=$?

env -u CUDA_DISABLE_PTX_JIT \
  CUDA_FORCE_PTX_JIT=1 CUDA_CACHE_DISABLE=1 \
  ./check > ptx-jit.log 2>&1
ptx_rc=$?

printf 'native=%s\nptx_jit=%s\n' \
  "$native_rc" "$ptx_rc" > execution-status.txt

cat execution-status.txt native.log ptx-jit.log
test "$native_rc" -eq 0 && test "$ptx_rc" -eq 0

두 로그 모두에 PASS가 출력되어야 하며, 종료 코드는 모두 0이어야 합니다. 테스트용 환경 변수는 해당 명령에만 적용합니다. 전역 환경 설정에 추가하여 업무 프로세스에 JIT를 상시 강제하지 마세요.

네이티브 실행은 실패하고 PTX 실행은 성공한다면 네이티브 대상과 최종 링크 단계를 먼저 점검합니다. 네이티브 실행은 성공하고 PTX 실행은 실패한다면 PTX가 포함되어 있는지, 드라이버가 해당 버전을 지원하는지 확인합니다. 둘 다 실패한다면 구체적인 오류를 바탕으로 툴체인, 런타임 라이브러리, 디바이스 상태를 점검해야 합니다. 원인을 파악하는 대신 컴파일 옵션만 반복해서 바꾸어서는 안 됩니다.

5. 실험 결과를 릴리스 승인 기준으로 만들기

이 소형 프로그램이 통과했다는 것은 기본 빌드 환경과 이 커널이 현재 환경에서 동작한다는 사실만 입증합니다. 실제 업무 프로그램은 확장 라이브러리를 로드하거나, 입력에 따라 다른 커널을 선택하거나, 런타임에 코드를 생성할 수도 있습니다. 따라서 실제 진입점을 사용하는 테스트를 추가하여 주요 데이터 유형과 분기를 검증해야 합니다.

릴리스마다 다음 네 가지 증빙을 보관하는 것이 좋습니다.

• 호스트 아키텍처, 드라이버 출력, 툴킷 버전, 실제 컴파일 명령.
• GPU 연산 능력과 최종 파일의 SASS 및 PTX 검사 결과.
• 두 실행 경로의 종료 코드, 로그, 업무 처리 결과 대조 기록.
• 소스 코드 버전, 최종 파일의 SHA-256, 동적 의존성 목록.

sha256sum ./check로 이번 검증 대상을 특정할 수 있습니다. 이후 다시 링크했다면 파일 이름이 같더라도 재검증해야 합니다. 컨테이너를 사용한다면 실제 배포 이미지 안에서 툴체인과 최종 파일을 검사하고, 실제 실행 환경에서 실행 검증을 완료해야 합니다.

PTX는 호환성을 위한 경로이지 모든 드라이버에서의 동작을 보장하는 수단은 아닙니다. 한 번 성공했다고 다른 디바이스에서도 동작한다고 결론 내려서는 더더욱 안 됩니다. 릴리스 승인 조건은 “현재 드라이버, 현재 산출물, 검증한 진입점에서 통과”로 명확히 규정해야 합니다. 아키텍처나 툴체인이 바뀌면 이전 로그를 재사용하지 말고 증빙을 새로 확보해야 합니다.

자주 묻는 질문