DGX Spark에서 CUDA 범위 초과 접근을 찾는 방법

CUDA 커널은 실행됐는데 결과를 호스트로 가져오는 단계에서야 오류가 발생하고, 같은 프로그램도 입력 크기를 줄이면 정상 동작하는 경우가 있습니다. 이런 증상이라면 드라이버부터 바꾸기보다 디바이스 메모리 접근을 먼저 점검해야 합니다. DGX Spark에서는 작은 독립 프로그램으로 범위 초과 접근을 의도적으로 일으키고, 문제가 발생한 소스 줄을 찾고, 인덱스를 수정한 뒤 오류 감지 기능이 꺼지지 않았는지까지 확인할 수 있습니다. 이 과정에는 모델 가중치가 필요하지 않으며 운영 중인 서비스에도 손대지 않습니다.

1. 환경을 고정하고 테스트 범위 정하기

TcpNode의 DGX Spark는 NVIDIA GB10 Grace Blackwell, ARM64 아키텍처, NVIDIA DGX OS를 사용하며 128GB LPDDR5x 통합 메모리를 탑재합니다. 통합 메모리를 CPU와 GPU가 함께 사용한다고 해서 임의의 포인터로 할당 경계를 넘어 접근해도 되는 것은 아닙니다. 여기서는 cudaMalloc으로 디바이스 메모리를 할당해 유효한 주소 범위를 명확히 합니다.

결제 후 콘솔에서 접속 정보를 확인하고 SSH로 로그인합니다. 테스트는 별도의 작업 디렉터리에서 일반 사용자 권한으로 실행해야 하며, 요청을 처리 중인 추론 프로세스에 오류를 주입해서는 안 됩니다.

mkdir -p "$HOME/cuda-access-check"
cd "$HOME/cuda-access-check"
uname -m
nvidia-smi
command -v nvcc
command -v compute-sanitizer
nvcc --version
compute-sanitizer --version

아키텍처는 aarch64로 표시되어야 합니다. 도구 버전과 드라이버 정보를 기록하고, 이후 재검증에도 같은 환경을 사용합니다. CUDA 스택이 사전 설치되어 있어도 모든 도구가 현재 PATH에 등록되어 있다는 뜻은 아닙니다. 명령을 찾을 수 없다면 설치된 툴체인과 설치 경로부터 확인하고, 드라이버 전체를 곧바로 교체하지 마세요.

범위 초과 접근 예제는 폐기해도 되는 독립 프로세스에서만 실행해야 합니다. 잘못된 메모리 접근이 발생하면 해당 프로세스를 종료하고 새로 실행하세요. 같은 CUDA 컨텍스트에서 얻은 후속 결과를 유효한 검증 근거로 삼아서는 안 됩니다.

2. 변수 하나만 바꾸는 범위 초과 접근 예제 만들기

다음 내용을 repro.cu로 저장합니다. 정수 하나만 할당하므로 유효한 인덱스는 0뿐입니다. 컴파일할 때 TEST_INDEX로 쓰기 위치를 바꿉니다. 오류 버전과 수정 버전이 같은 소스를 사용하므로 입력, 할당 크기, 실행 경로를 동시에 바꾸지 않고 비교할 수 있습니다.

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

static void checked(cudaError_t status) {
    if (status != cudaSuccess) {
        std::fprintf(stderr, "%s\n", cudaGetErrorString(status));
        std::exit(EXIT_FAILURE);
    }
}

__global__ void write_value(int* data) {
    data[TEST_INDEX] = 7;
}

int main() {
    int* data = nullptr;
    int result = 0;
    checked(cudaMalloc(reinterpret_cast<void**>(&data), sizeof(int)));
    checked(cudaMemset(data, 0, sizeof(int)));
    write_value<<<1, 1>>>(data);
    checked(cudaGetLastError());
    checked(cudaDeviceSynchronize());
    checked(cudaMemcpy(&result, data, sizeof(int), cudaMemcpyDeviceToHost));
    checked(cudaFree(data));
    std::printf("result=%d\n", result);
    return result == 7 ? EXIT_SUCCESS : EXIT_FAILURE;
}

두 프로그램을 각각 컴파일합니다.

nvcc -arch=native -lineinfo -DTEST_INDEX=1 repro.cu -o repro-bad
nvcc -arch=native -lineinfo -DTEST_INDEX=0 repro.cu -o repro-good
sha256sum repro.cu repro-bad repro-good > build.sha256

-arch=native는 현재 인식되는 GPU를 대상으로 컴파일하는 옵션입니다. 현재 장비에서 검증할 때 적합하며, 여러 디바이스에서 공통으로 사용할 바이너리를 만드는 용도로는 쓰지 않습니다. 컴파일러가 현재 디바이스를 지원하지 않는다면 툴체인 호환성부터 확인하고, 다른 아키텍처용 옵션을 임의로 적용하지 마세요. -lineinfo는 디바이스 소스의 줄 정보를 보존합니다. 처음 오류 위치를 찾을 때부터 실행 특성을 크게 바꾸는 전체 디바이스 디버그 모드를 켤 필요는 없습니다.

커널 실행 요청과 동기화를 모두 확인해야 하는 이유

cudaGetLastError는 커널 실행 요청과 관련된 오류를 확인하지만, 커널은 보통 비동기로 실행됩니다. cudaDeviceSynchronize는 앞서 제출된 디바이스 작업이 끝날 때까지 기다리므로 실행 중 발생한 오류가 이 시점에 반환될 수 있습니다. 두 호출을 각각 확인해야 하며, 프로그램이 결과를 출력했는지만 봐서는 안 됩니다.

3. 감지 결과를 판정 가능한 종료 코드로 전달하기

먼저 오류 버전을 실행하고 표준 출력과 감지 로그를 따로 저장합니다.

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-bad > bad.stdout 2> bad.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat bad.stdout bad.stderr

잘못된 전역 메모리 쓰기가 보고되고, write_value 안의 쓰기 연산이 문제 지점으로 표시되어야 합니다. 보고서에는 일반적으로 접근 크기, 스레드 위치, 할당 영역과 접근 주소의 관계도 포함됩니다. 표시 형식은 도구 버전에 따라 달라지므로 특정 문구 하나만을 판정 기준으로 삼아서는 안 됩니다.

--error-exitcode 86은 감지된 오류가 자동화 작업의 성공·실패 판정에 반영되도록 하는 옵션입니다. 다만 테스트 대상 프로그램 자체가 이미 실패했다면 프로그램이 반환한 다른 0이 아닌 값이 나올 수도 있습니다. 따라서 통과 기준은 “종료 코드가 0이 아니고 보고서가 의도한 범위 초과 접근을 가리키는가”이지, 모든 실패에서 반드시 86을 반환하는가는 아닙니다.

명령을 일반적인 로그 파이프라인에 바로 연결한 뒤 마지막 로그 명령의 상태만 확인하지 마세요. 위의 리디렉션과 조건 분기는 실제 실행 결과를 보존하며, 실패 즉시 종료하도록 설정된 스크립트에서도 사용할 수 있습니다.

4. 증상을 숨기지 말고 주소 범위에 맞게 수정하기

이 예제는 정수 하나만 할당했는데 쓰기 인덱스가 1이므로 유효 범위를 벗어납니다. 실제 프로그램에서는 보고서에 나온 주소를 해당 할당 영역과 연결한 다음 인덱스 계산, 요소 타입, 스트라이드, 배치 오프셋을 확인해야 합니다.

흔한 원인으로는 할당 인수에는 바이트 수를 쓰면서 인덱스에서는 요소 수로 해석하는 경우, 마지막 스레드 블록을 채울 만큼 요소가 남지 않는 경우, 2차원 배열의 행 스트라이드가 논리적 너비와 다른 경우, 해제된 포인터를 다른 작업이 다시 사용하는 경우 등이 있습니다. 할당 크기를 늘리면 오류가 일시적으로 드러나지 않을 수는 있지만, 알고리즘이 올바르다는 증거가 되지는 않습니다.

1차원 작업에서는 일반적으로 접근 전에 index < count를 확인하고, 인덱스 계산에 사용하는 정수 타입이 곱셈 결과를 담을 수 있는지도 확인해야 합니다. 공유 메모리로 스레드가 협업하는 커널에서는 일부 스레드만 임의로 조기 반환시킨 뒤 나머지 스레드가 블록 단위 배리어에 진입하게 해서는 안 됩니다. 새로운 동기화 문제가 생길 수 있기 때문입니다.

대규모 워크플로에서만 오류가 발생한다면 의심되는 단계 뒤에 동기화 검사를 임시로 추가해 구간별로 범위를 좁힐 수 있습니다. 검사 지점은 한 번에 하나씩만 추가하고 최초 오류를 기록하세요. 오류가 발생한 뒤에도 작업을 계속 제출해 대량의 후속 오류를 만들지 않도록 합니다.

5. 정상·실패 테스트로 수정 검증 마무리하기

수정 버전을 실행합니다.

if compute-sanitizer --tool memcheck --error-exitcode 86 \
    ./repro-good > good.stdout 2> good.stderr; then
    status=0
else
    status=$?
fi
printf 'exit=%s\n' "$status"
cat good.stdout good.stderr

이 예제에서는 종료 코드가 0이고, 감지 보고서에 메모리 접근 오류가 없으며, 출력이 result=7이라는 세 조건을 모두 만족해야 합니다. “오류가 나오지 않았다”는 사실만으로는 부족합니다. 커널이 실행되지 않았거나 결과가 실제로 기록되지 않았어도 겉으로는 테스트를 통과할 수 있습니다.

이어서 오류 버전을 다시 실행하고 여전히 실패하는지 확인합니다. 이 실패 검증을 통해 감지 명령이 삭제된 경우, 테스트에서 실수로 오래된 바이너리를 사용한 경우, 파이프라인이 실패 상태를 누락한 경우 등을 찾을 수 있습니다. 실제 커널에서는 최소 유효 입력, 블록 경계에 정확히 맞는 입력, 마지막 블록을 채우지 못하는 입력도 검증해야 합니다. 인터페이스가 빈 입력을 허용한다면 처리 방식도 명확히 해야 합니다.

마지막으로 소스, 컴파일 명령, 버전 정보, 체크섬, 두 실행의 로그를 보관합니다. 감지 도구는 실행 오버헤드를 바꾸므로 검사 중 측정한 시간을 그대로 성능 평가에 사용할 수 없습니다. 또한 memcheck를 통과했다고 해서 데이터 경쟁과 동기화 문제까지 배제된 것은 아닙니다. 검증을 제대로 마무리하려면 성공 화면 하나만 남길 것이 아니라, 메모리 접근의 유효성, 계산 결과, 실패 상태의 전달을 각각 다시 확인할 수 있는 근거를 갖춰야 합니다.

자주 묻는 질문