DGX Spark에서 CUDA 스트림 의존성과 버퍼 수명 검증하기

두 CUDA 스트림이 같은 디바이스 버퍼를 공유하면 여러 번 연속으로 정상 동작하다가도 입력 크기를 바꾸는 순간 이전 값이 나타날 수 있습니다. 흔한 원인은 연산식이 아니라, 생산자 스트림이 쓰기를 끝내기 전에 소비자 스트림이 읽기를 시작하는 데 있습니다. 이 글에서는 DGX Spark에서 두 스트림을 사용하는 최소 실험을 구성하고, 작업 제출 순서와 실행 의존성, 버퍼 재사용 시점을 각각 검증합니다. 한 번 실행에 성공했다는 이유만으로 올바른 프로그램이라고 판단하지 않습니다.

1. 스트림을 만들기 전에 데이터 의존성부터 정리하기

DGX Spark는 NVIDIA GB10 Grace Blackwell을 사용하며 128GB LPDDR5x 통합 메모리를 제공합니다. CPU와 GPU는 같은 메모리 풀을 공유하지만, 메모리가 물리적으로 통합되어 있다고 해서 애플리케이션의 스트림 간 실행 순서까지 자동으로 보장되지는 않습니다. 명시적 복사, 커널 실행, 호스트 접근은 여전히 각각의 완료 조건을 따라야 합니다.

이 실험에서는 처리 경계를 쉽게 확인할 수 있도록 디바이스 메모리를 명시적으로 할당하고 페이지 고정 호스트 메모리를 사용합니다. 생산자 스트림은 “입력을 디바이스로 복사한 뒤 각 요소에 2를 곱하는” 작업을, 소비자 스트림은 “각 요소에 1을 더한 뒤 결과를 호스트로 복사하는” 작업을 수행합니다. 전체 의존성은 다음과 같습니다.

입력 준비 완료 → 복사 완료 → 2배 연산 완료 → 1 더하기 연산 시작 → 호스트로 전송 완료 → 호스트에서 검사.

같은 스트림 안의 작업은 스트림 순서대로 실행되지만, 두 스트림 사이는 명시적으로 연결해야 합니다. 생산자 스트림에서 2배 연산 커널 뒤에 이벤트를 기록하고, 소비자 스트림에서는 해당 이벤트를 기다린 다음 1 더하기 커널을 제출해야 합니다.

호스트가 생산자 스트림을 먼저 호출하고 소비자 스트림을 나중에 호출했다는 것은 제출 순서만 의미합니다. 스트림 간 데이터 의존성에는 명확한 동기화 근거가 있어야 하며, 부하가 적을 때 우연히 성립하는 실행 순서에 의존해서는 안 됩니다.

2. 기본 스트림에 의존하지 않는 최소 프로그램 만들기

TcpNode에서 결제한 뒤 콘솔에서 접속 정보를 확인하고 SSH로 NVIDIA DGX OS에 접속합니다. 먼저 uname -m, nvidia-smi, nvcc --version을 실행해 호스트 아키텍처, 디바이스 및 드라이버 정보, 컴파일 도구 체인 정보를 각각 저장합니다. 아키텍처는 ARM64에 해당하는 aarch64여야 합니다. 컴파일러 버전을 드라이버 출력으로 추정해서는 안 됩니다.

다음 내용을 stream_contract.cu로 저장합니다. 두 스트림 모두 논블로킹 플래그를 사용해, 기존 기본 스트림의 암묵적 동기화가 누락된 의존성을 가리지 않도록 합니다. 이벤트는 순서 제어에만 사용하며 시간을 측정하지 않습니다.

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

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

__global__ void scale(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] *= 2.0f;
}

__global__ void add(float* data, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[i] += 1.0f;
}

int main() {
    constexpr int count = 4096;
    const size_t bytes = count * sizeof(float);
    float *input = nullptr, *output = nullptr, *device = nullptr;
    cudaStream_t producer, consumer;
    cudaEvent_t ready;

    check(cudaMallocHost(reinterpret_cast<void**>(&input), bytes));
    check(cudaMallocHost(reinterpret_cast<void**>(&output), bytes));
    check(cudaMalloc(reinterpret_cast<void**>(&device), bytes));
    check(cudaStreamCreateWithFlags(&producer, cudaStreamNonBlocking));
    check(cudaStreamCreateWithFlags(&consumer, cudaStreamNonBlocking));
    check(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));

    for (int i = 0; i < count; ++i) {
        input[i] = static_cast<float>(i);
        output[i] = -1.0f;
    }

    check(cudaMemcpyAsync(device, input, bytes,
                          cudaMemcpyHostToDevice, producer));
    scale<<<(count + 255) / 256, 256, 0, producer>>>(device, count);
    check(cudaGetLastError());
    check(cudaEventRecord(ready, producer));

    check(cudaStreamWaitEvent(consumer, ready, 0));
    add<<<(count + 255) / 256, 256, 0, consumer>>>(device, count);
    check(cudaGetLastError());
    check(cudaMemcpyAsync(output, device, bytes,
                          cudaMemcpyDeviceToHost, consumer));
    check(cudaStreamSynchronize(consumer));

    int failures = 0;
    for (int i = 0; i < count; ++i) {
        float expected = 2.0f * static_cast<float>(i) + 1.0f;
        if (output[i] != expected) ++failures;
    }
    std::printf("elements=%d mismatches=%d\n", count, failures);

    check(cudaEventDestroy(ready));
    check(cudaStreamDestroy(consumer));
    check(cudaStreamDestroy(producer));
    check(cudaFree(device));
    check(cudaFreeHost(output));
    check(cudaFreeHost(input));
    return failures == 0 ? 0 : 2;
}

3. 출력과 종료 코드를 함께 검증하기

같은 디렉터리에서 컴파일하고 실행합니다.

nvcc -O2 stream_contract.cu -o stream_contract
./stream_contract
printf 'exit=%s\n' "$?"

통과 조건은 출력이 elements=4096 mismatches=0이고 종료 코드가 0인 것입니다. API 호출이나 동기화에 실패하면 1을, 계산값이 일치하지 않으면 2를 반환합니다. 프로세스가 비정상 종료되지 않았는지만 확인하면, 계산은 끝났지만 결과가 잘못된 경우를 놓칠 수 있습니다.

예제의 값은 모두 단정밀도 부동소수점으로 정확하게 표현할 수 있는 정수 범위에 있으므로 동등 비교를 직접 사용합니다. 실제 부동소수점 모델에 이 기준을 그대로 적용해서는 안 됩니다. 알고리즘에 맞춰 절대 오차, 상대 오차, 유한하지 않은 값의 처리 규칙을 정해야 합니다.

이어서 요소 수를 각각 1, 257, 4097로 바꾸고 다시 컴파일해 실행합니다. 이 크기들은 단일 요소, 스레드 블록의 끝부분, 여러 블록에 걸친 경계를 확인하기 위한 것입니다. 성능 데이터가 아니라 테스트 입력입니다. 성공 화면 한 장만 남기지 말고, 매번 소스 파일과 컴파일 명령, 도구 체인 정보, 실행 출력을 함께 저장하세요.

실패 감지 경로가 실제로 동작하는지 확인하기

기댓값 식 끝의 1.0f를 잠시 2.0f로 바꿉니다. 프로그램이 불일치를 보고하고 2를 반환하는지 확인한 다음 원래 식으로 복원합니다. 이렇게 하면 결과 검사 기능이 제대로 동작하는지 검증할 수 있습니다. “대기를 삭제하면 반드시 오류가 나야 한다”는 조건으로 부정 테스트를 해서는 안 됩니다. 경쟁 상태가 드러나는지는 스케줄링에 달려 있으므로, 의존성이 빠진 프로그램도 일시적으로 테스트를 통과할 수 있습니다.

4. 이벤트 경계를 메모리 수명에 연결하기

일반적으로 cudaMemcpyAsync가 반환되었다는 것은 호출이 제출되었다는 의미일 뿐입니다. 이를 근거로 입력을 덮어써서는 안 됩니다. 예제에서는 복사를 제출한 시점부터 소비자 스트림의 동기화가 끝날 때까지 입력을 변경하지 않습니다. 출력도 호스트로 전송하는 스트림의 동기화가 끝난 뒤에만 읽습니다. 보수적인 전략이지만 경계가 명확하므로, 정확성을 검증하는 초기 기준으로 적합합니다.

페이지 고정 메모리는 비동기 전송 조건을 충족하는 데 도움이 되지만 완료 알림은 아닙니다. 디바이스 버퍼 역시 다른 스트림이 접근 중일 때 미리 해제해서는 안 됩니다. 또한 128GB 통합 메모리가 있다고 해서 대량의 호스트 메모리를 모두 페이지 고정해야 한다는 뜻은 아닙니다. 배치 용량은 실제 작업 집합을 기준으로 제한해야 합니다.

생산자 이벤트가 보장하는 범위는 그 이벤트 앞에 있는 작업뿐이며, 소비자 스트림의 1 더하기 연산과 호스트로의 전송은 포함하지 않습니다. 따라서 생산자 이벤트가 완료되면 생산 단계가 끝났다는 것은 확인할 수 있지만, 다음 배치가 디바이스 버퍼를 덮어써도 된다는 뜻은 아닙니다. 파이프라인으로 확장할 때는 소비자 스트림에서 호스트로의 전송 뒤에 완료 이벤트를 추가로 기록하고, 이를 기준으로 버퍼 전체의 재사용을 제어해야 합니다.

병렬 처리에 사용하는 각 배치 슬롯에는 독립적인 입력, 출력, 디바이스 버퍼, 완료 상태가 있어야 합니다. 먼저 단일 슬롯으로 검증한 다음 슬롯 수를 늘리세요. 스트림 수만 늘리고 같은 포인터를 아무런 보호 없이 계속 공유해서는 안 됩니다.

5. 문제를 추적하며 동기화 범위를 단계적으로 좁히기

결과가 간헐적으로 틀린다면 먼저 대기가 소비자 커널보다 앞에 있는지, 이벤트 기록이 생산자 커널보다 뒤에 있는지, 기록과 대기에 같은 이벤트를 사용하는지 확인합니다. 이벤트 기록 위치가 한 줄만 앞당겨져도 “데이터 쓰기 완료”라는 조건이 잘못해서 “복사 완료”로 축소될 수 있습니다.

다음으로 각 커널 실행 직후의 오류와 최종 동기화의 반환값을 확인합니다. 전자는 일부 실행 시작 문제를 감지하고, 후자는 비동기 실행 중 발생한 오류를 확인하는 역할을 합니다. 둘 중 하나로 다른 하나를 대체할 수는 없습니다. 실행 시작 시점의 반환값만 확인하고 완료를 기다리지 않으면, 오류가 훨씬 뒤의 API 호출에서야 드러날 수 있습니다.

진단용 사본에서는 생산 단계에 스트림 동기화를 임시로 추가해 문제가 스트림 간 순서와 관련 있는지 판단할 수 있습니다. 오류가 사라진다면 디바이스 전체 동기화로 마무리하지 말고, 의존성과 버퍼 소유권을 계속 확인해야 합니다. 전역 동기화는 스케줄링을 바꾸며, 원래 있던 수명 관리 문제를 감출 수 있습니다.

최종적으로 남겨야 할 것은 명확한 생산자 이벤트, 소비자 대기, 버퍼 재사용을 제어하는 완료 이벤트, 그리고 반복 실행 가능한 출력 검사입니다. 의존 관계와 메모리 소유권을 모두 명확히 설명할 수 있어야 멀티 스트림 프로그램의 동시성을 더 확장할 기반이 갖춰집니다.

자주 묻는 질문