CUDA C/C++ 기초

References

Fundamentals of Accelerated Computing with CUDA C/C++ (NVIDIA Online Training)
Asynchronous Streaming, and Visual Profiling for Accelerated Applications with CUDA C/C++

Concurrent CUDA Streams
Manual Device Memory Allocation and Copying
Using Streams to Data Transfers

지난 두 포스팅을 통해서 기본적인 CUDA와 Unified Memory에 대해서 살펴봤습니다. 또한, Nsight System의 프로파일링을 통해 정량적으로 성능을 평가했습니다.

이번 포스팅에서는 CUDA의 동시성(concurrent)에 대해서 살펴보도록 하겠습니다.

Concurrent CUDA Streams

CUDA 프로그래밍에서 스트림(stream)은 순서대로 실행되는 일련의 명령(instructions)입니다. CUDA에서 커널과 일부 메모리 전송(transfer)은 CUDA 스트림 내에서 발생합니다. 이전 포스팅까지 명시적으로 CUDA 스트림과 상호작용하지는 않았습니다만, 기본적으로 CUDA 커널은 디폴트 스트림(default stream)에서 실행됩니다.

CUDA에서는 디폴트 스트림 외에 다른 스트림을 생성하고 활용할 수 있습니다. 이를 통해서 여러 커널들을 다른 스트림에서 동시에 실행하여 여러 작업을 수행할 수 있습니다. 다중 스트림을 사용하면 병렬화 계층을 추가할 수 있고, 더 최적화할 수 있는 기회가 주어집니다.

어떤 하나의 단일 스트림 내에서의 커널은 반드시 순서대로 실행됩니다. 하지만, 다른 스트림(non-default)에 있는 각각의 커널들을 동시에 수행될 수 있습니다.

디폴트 스트림은 조금 특별한데, 디폴트 스트림은 다른 모든 스트림의 모든 커널을 블락(block) 합니다.

따라서, CUDA 스트림은 다음과 같은 규칙을 가지고 있습니다.

주어진 스트림 내에서의 명령(operations)은 순차적으로 발생한다.
non-default인 다른 스트림에서의 명령은 특정한 순서로 동작한다고 보장되지 않는다. 즉, 다른 스트림에서의 명령의 순서는 고정되어 있지 않다.
디폴트(default) 스트림은 실행되기 전에 다른 모든 스트림이 완료될 때까지 기다리며, 디폴트 스트림의 실행이 완료될 때까지 다른 스트림이 실행되지 않도록 블락한다.

Creating, Utilizing, and Destroying Non-Default CUDA Streams

아래의 예제 코드는 non-default 스트림을 만들고, 활용하고, 제거하는 방법을 보여줍니다. 디폴트가 아닌 스트림에서 커널을 실행하려면 execution configuration의 4번째 인수로 스트림이 전달되어야 합니다.

cudaStream_t stream;       // CUDA Streams are of type 'cudaStream_t'
cudaStreamCreate(&stream); // Note that a pointer must be passed to 'cudaCreateStream'

someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>();

cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to 'cudaDestroyStream'

참고로 execution configuration의 3번째 인수는 각 블록에서 사용할 shared memory의 바이트 수를 전달합니다. 여기서는 0으로 전달하며, 아래의 예제에서 shared memory를 사용하지는 않지만 4번째 인수를 전달하기 위해서 3번째 인수를 표기합니다.

#include <stdio.h>

__global__
void printNumber(int number)
{
    printf("%d\n", number);
}

int main()
{
    for (int i = 0; i < 5; i++) {
        printNumber<<<1, 1>>>(i);
    }

    cudaDeviceSynchronize();
}

위 예제 코드를 컴파일 후, 실행해보도록 하겠습니다. printNumber 커널이 디폴트 스트림에서 실행되기 때문에 0부터 4가 순차적으로 출력되는 것을 확인할 수 있습니다.

5개의 커널이 모두 동일한 스트림에서 실행되었기 때문에 5개의 커널이 순차적으로 실행되었습니다. 또한 디폴트 스트림은 블로킹하기 때문에 커널의 각 실행이 완료될 때까지 다음 커널의 실행이 대기하게 됩니다.

이제 코드를 아래와 같이 변경하여, 각 커널의 실행이 디폴트 스트림이 아닌 각각의 자체 스트림에서 발생되도록 합니다.

#include <stdio.h>

__global__
void printNumber(int number)
{
    printf("%d\n", number);
}

int main()
{
    for (int i = 0; i < 5; i++) {
        cudaStream_t stream;
        cudaStreamCreate(&stream);
        printNumber<<<1, 1, 0, stream>>>(i);
        cudaStreamDestroy(stream);
    }

    cudaDeviceSynchronize();
}

아마 위와 같이 그대로 순차적으로 출력될 것 입니다. 하지만, nsys profile을 통해서 세부적인 프로파일링 결과를 살펴보면 다르다는 것을 확인할 수 있습니다.

이전 코드와 현재 코드의 프로파일링 결과에서 cudaLaunchKernel API의 Total Time과 Num Calls에 주목해봅시다. printNumber커널을 다섯 번 호출하기 때문에 이 API의 호출도 5인 것을 확인할 수 있습니다. 여기까지는 동일하지만, Total Time을 보시면 차이가 크다는 것을 볼 수 있습니다. 디폴트 스트림을 사용한 경우에는 모든 커널이 순차적으로 실행하기 때문에 cudaLaunchKernel의 총 실행 시간이 238ms라는 것을 볼 수 있습니다. 반면에 각각 자체 스트림을 사용한 경우에는 0.07ms로 매우 감소한 것을 확인할 수 있습니다.

이 결과는 Nsight System을 UI로 실행하여 리포트를 열어보면, 그 차이점을 확실하게 확인할 수 있습니다.

The result that kernels run in default-stream.

The result that kernels run in each stream(non-default)

순차적으로 실행되던 printNumber 커널들이 각각의 non-default 스트림에서 실행하면 서로 중첩되어 동시에 실행되는 것을 시각적으로 확인할 수 있습니다.

아래 코드는 지난 포스팅에서 벡터 덧셈 어플리케이션에 prefetch 테크닉을 적용한 코드입니다.

#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

__global__
void initWith(float num, float *a, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride)
    {
        a[i] = num;
    }
}

__global__
void addVectorsInto(float* result, float* a, float* b, const int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride) {
        result[i] = a[i] + b[i];
    }
}

void checkElementsAre(float target, float* array, const int N)
{
    for (int i = 0; i < N; i++) {
        if (array[i] != target) {
            printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
            exit(1);
        }
    }
    printf("SUCCESS! All values added correctly.\n");
}

int main()
{
    const int N = 2 << 20;
    size_t size = N * sizeof(float);

    int deviceId;
    checkCuda(cudaGetDevice(&deviceId));

    cudaDeviceProp props;
    checkCuda(cudaGetDeviceProperties(&props, deviceId));

    float *a, *b, *c;

    checkCuda(cudaMallocManaged(&a, size));
    checkCuda(cudaMallocManaged(&b, size));
    checkCuda(cudaMallocManaged(&c, size));

    checkCuda(cudaMemPrefetchAsync(a, size, deviceId));
    checkCuda(cudaMemPrefetchAsync(b, size, deviceId));
    checkCuda(cudaMemPrefetchAsync(c, size, deviceId));

    size_t threadsPerBlock = props.maxThreadsPerBlock;
    size_t numberOfBlocks = props.multiProcessorCount;

    initWith<<<numberOfBlocks, threadsPerBlock>>>(3, a, N);
    initWith<<<numberOfBlocks, threadsPerBlock>>>(4, b, N);
    initWith<<<numberOfBlocks, threadsPerBlock>>>(0, c, N);

    addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);

    checkCuda(cudaGetLastError());
    checkCuda(cudaDeviceSynchronize());

    checkCuda(cudaMemPrefetchAsync(c, size, cudaCpuDeviceId));
    
    checkElementsAre(7, c, N);

    checkCuda(cudaFree(a));
    checkCuda(cudaFree(b));
    checkCuda(cudaFree(c));
}

위 코드에서 조금 더 최적화할 여지가 있는데, 바로 초기화 커널인 initWith이 위 코드에서는 디폴트 스트림에서 실행되기 때문에 순차적으로 실행됩니다. 초기화 커널의 실행은 서로 영향을 미치지 않기 때문에 병렬로 실행되어도 무관합니다. 따라서, 각각의 스트림을 할당하여 동시에 커널이 시작되도록 수정할 수 있습니다.

#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

__global__
void initWith(float num, float *a, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride)
    {
        a[i] = num;
    }
}

__global__
void addVectorsInto(float* result, float* a, float* b, const int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride) {
        result[i] = a[i] + b[i];
    }
}

void checkElementsAre(float target, float* array, const int N)
{
    for (int i = 0; i < N; i++) {
        if (array[i] != target) {
            printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
            exit(1);
        }
    }
    printf("SUCCESS! All values added correctly.\n");
}

int main()
{
    const int N = 2 << 20;
    size_t size = N * sizeof(float);

    int deviceId;
    checkCuda(cudaGetDevice(&deviceId));

    cudaDeviceProp props;
    checkCuda(cudaGetDeviceProperties(&props, deviceId));

    float *a, *b, *c;

    checkCuda(cudaMallocManaged(&a, size));
    checkCuda(cudaMallocManaged(&b, size));
    checkCuda(cudaMallocManaged(&c, size));

    checkCuda(cudaMemPrefetchAsync(a, size, deviceId));
    checkCuda(cudaMemPrefetchAsync(b, size, deviceId));
    checkCuda(cudaMemPrefetchAsync(c, size, deviceId));

    size_t threadsPerBlock = props.maxThreadsPerBlock;
    size_t numberOfBlocks = props.multiProcessorCount;

    cudaStream_t stream1, stream2, stream3;
    checkCuda(cudaStreamCreate(&stream1));
    checkCuda(cudaStreamCreate(&stream2));
    checkCuda(cudaStreamCreate(&stream3));

    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);

    addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);

    checkCuda(cudaGetLastError());
    checkCuda(cudaDeviceSynchronize());

    checkCuda(cudaMemPrefetchAsync(c, size, cudaCpuDeviceId));
    
    checkElementsAre(7, c, N);

    checkCuda(cudaStreamDestroy(stream1));
    checkCuda(cudaStreamDestroy(stream2));
    checkCuda(cudaStreamDestroy(stream3));

    checkCuda(cudaFree(a));
    checkCuda(cudaFree(b));
    checkCuda(cudaFree(c));
}

default 스트림만 사용했을 때는 cudaLaunchKernel의 총 실행 시간이 약 0.057ms였는데, non-default 스트림을 사용하여 초기화를 수행했을 때에는 총 실행 시간이 0.050ms로 감소하였습니다.

Manual Device Memory Allocation and Copying

지금까지 사용했던 cudaMallocManaged와 cudaMemPrefetchAsync는 성능이 꽤 좋고 메모리 마이그레이션을 간단하게 수행할 수 있도록 하지만, 아마 상용 어플리케이션에서는 수동으로 메모리를 할당하는 방법이 더 많이 사용됩니다. 특히, 데이터가 device 또는 host에서만 액세스되기 때문에 자동으로 마이그레이션하는 기능이 필요없어서 이러한 오버헤드를 제거해야하는 경우가 특히 그렇습니다.

또한, UM을 사용하지 않고 수동으로 메모리를 관리하면, non-default 스트림을 사용하여 연산과 데이터 전달을 오버랩할 수 있습니다. 따라서, UM을 사용하는 것이 아닌 수동으로 CPU/GPU에 메모리를 할당하는 방법도 중요합니다.

수동으로 메모리를 관리하기 위해 아래와 같은 CUDA API 함수들이 있습니다.

cudaMalloc을 사용하면 활성화된 GPU에 직접 메모리를 할당합니다. 이 API를 사용하면 모든 GPU의 페이지 폴트를 방지합니다. 대신, 이 API가 반환하는 포인터는 host 코드에서 액세스가 불가능합니다.
cudaMallocHost는 CPU에 직접 메모리를 할당합니다. 또한, 메모리를 고정(pin)하거나 page-lock을 통해 GPU에서 메모리를 비동기로 복사할 수 있도록 합니다. 이렇게 할당된 메모리를 Pinned Memory라고 하며, 이렇게 할당된 메모리가 너무 많으면 CPU의 성능을 하락시킬 수 있으므로 명확한 의도를 가지고 사용해야 합니다. 이렇게 할당된 메모리는 cudaFreeHost API를 통해 해제할 수 있습니다.
cudaMemcpy는 메모리를 복사(전달이 아님)합니다. host에서 device 또는 그 역방향으로도 복사할 수 있습니다.

필요하시면, Pinned Memory에 관련된 내용은 아래 포스팅에서 참조바랍니다 !

Pinned Memory

References Professional CUDA C Programming Contents Pinned Memory Pinned Memory 할당된 Host 메모리는 기본적으로 pageable합니다. 즉, OS에 의해 host의 가상 메모리에서 다른 물리 메모리로 데이터를 이동..

junstar92.tistory.com

Manual Device Memory Management Example

아래 예제 코드는 위에서 언급한 CUDA API를 통해 어떻게 메모리를 할당하고 복사하고 해제하는 지 보여줍니다.

int *host_a, *device_a;         // Define host-specific and device-specific arrays.
cudaMalloc(&device_a, size);    // 'device_a' is immediately available on the GPU.
cudaMallocHost(&host_a, size);  // 'host_a' is immediately available on CPU, and is page-locked, or pinned.

initializeOnHost(host_a, N);    // No CPU page faulting since memory is already allocated on the host.

// 'cudaMemcpy' takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.
cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);

kernel<<<block, threads, 0, someStream>>>(device_a, N);

// 'cudaMemcpy' can also copy data from device to host.
cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);

verifyOnHost(host_a, N);

cudaFree(device_a);
cudaFreeHost(host_a);  // Free pinned memory like this.

지금까지 예제로 사용했던 벡터 덧셈 어플리케이션 코드는 cudaMallocManaged를 사용하여 메모리를 관리했습니다. 이번에는 수동으로 device 메모리를 할당하고 복사하게 되면 성능이 어떻게 변하게 되는지 살펴보도록 하겠습니다.

// 15_vector-add-manual-alloc.cu
#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

__global__
void initWith(float num, float *a, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride)
    {
        a[i] = num;
    }
}

__global__
void addVectorsInto(float* result, float* a, float* b, const int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride) {
        result[i] = a[i] + b[i];
    }
}

void checkElementsAre(float target, float* array, const int N)
{
    for (int i = 0; i < N; i++) {
        if (array[i] != target) {
            printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
            exit(1);
        }
    }
    printf("SUCCESS! All values added correctly.\n");
}

int main()
{
    const int N = 2 << 20;
    size_t size = N * sizeof(float);

    int deviceId;
    checkCuda(cudaGetDevice(&deviceId));

    cudaDeviceProp props;
    checkCuda(cudaGetDeviceProperties(&props, deviceId));

    float *a, *b, *c, *h_c;

    checkCuda(cudaMalloc(&a, size));
    checkCuda(cudaMalloc(&b, size));
    checkCuda(cudaMalloc(&c, size));
    checkCuda(cudaMallocHost(&h_c, size));

    size_t threadsPerBlock = props.maxThreadsPerBlock;
    size_t numberOfBlocks = props.multiProcessorCount;

    cudaStream_t stream1, stream2, stream3;
    checkCuda(cudaStreamCreate(&stream1));
    checkCuda(cudaStreamCreate(&stream2));
    checkCuda(cudaStreamCreate(&stream3));

    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);

    addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);

    checkCuda(cudaMemcpy(h_c, c, size, cudaMemcpyDeviceToHost));
    
    checkElementsAre(7, h_c, N);

    checkCuda(cudaStreamDestroy(stream1));
    checkCuda(cudaStreamDestroy(stream2));
    checkCuda(cudaStreamDestroy(stream3));

    checkCuda(cudaFree(a));
    checkCuda(cudaFree(b));
    checkCuda(cudaFree(c));
    checkCuda(cudaFreeHost(h_c));
}

프로파일링 결과는 다음과 같습니다.

지난 포스팅에서 UM을 사용한 결과와 비교했을 때, 가장 큰 차이점은 메모리를 할당하는 데 약 10ms 정도 감소하고, Memory Operation의 총 실행 시간 또한 3ms 정도 감소했다는 것입니다. 대신, Pinned Memory를 할당하고 해제하는 데 시간이 약 5ms 정도 추가되고, cudaMemcpy에 의한 API 실행 시간이 1.5ms 정도 추가되었기 때문에 드라마틱한 차이는 발생하지 않은 것 같습니다. 제 생각에는 Unified Memory를 사용했지만 비동기 Prefetch 테크닉을 통해 최소한의 오버헤드만 발생했기 때문으로 추측됩니다.

참고로 이전 포스팅에서 Unified Memory를 사용하고, Prefetch 테크닉을 적용한 코드의 프로파일링 결과는 다음과 같습니다.

Using Streams to Data Transfers

cudaMemcpy 외에도 메모리를 복사할 수 있는 cudaMemcpyAsync라는 API가 제공됩니다. 이 API를 사용하면, host의 메모리가 고정되어 있는 한 host에서 device로, 또는 device에서 host로 메모리를 비동기로 복사할 수 있습니다.

커실의 실행과 유사하게, cudaMemcpyAsync 또한 기본적으로 host에 대해서만 비동기입니다. 기본적으로 default 스트림에서 실행되기 때문에 GPU에서 발생하는 다른 CUDA 명령을 블락합니다. 그러나 cudaMemcpyAsync 함수는 5번째 인수로 non-default 스트림을 옵션으로 받습니다. 만약 non-default 스트림이 전달되면, 다른 non-default 스트림에서 발생하는 CUDA 명령과 동시에 메모리를 전송할 수 있습니다.

일반적으로 유용하게 사용되는 패턴은 고정된(pinned) host 메모리, non-default 스트림에서의 비동기 메모리 복사, 그리고 non-default 스트림에서의 커널 실행을 조합하여 커널의 실행과 메모리 전송을 오버랩하는 것입니다.

아래 예제 코드는 커널에서 작업을 시작하기 전에 전체 메모리 복사가 완료될 때까지 기다리지 않고 필요한 데이터 일부가 복사되고 작업이 수행되며, 각 copy/work 작업은 default가 아닌 자체 스트림(non-default)에서 실행됩니다. 이 테크닉을 사용하면 다음 데이터 일부에 대한 메모리 전송이 수행되는 동안 현재 데이터 일부에 대한 작업을 수행할 수 있습니다. 이 테크닉을 사용할 때는 작업할 데이터 수에 대한 값과 오프셋 위치 계산에 주의해야 합니다.

int N = 2 << 24;
int size = N * sizeof(int);

int *host_array;
int *device_array;

cudaMallocHost(&host_array, size); // Pinned host memory allocation.
cudaMalloc(&device_array, size);   // Allocation directly on the active GPU device.

initializeData(host_array, N); // Assume this application needs to initialize on the host.

const int numberOfSegments = 4;               // This example demonstrates slicing the work into 4 segments.
int segmentN = N / numberOfSegments;          // A value for a segment's worth of `N` is needed.
size_t segmentSize = size / numberOfSegments; // A value for a segment's worth of `size` is needed.

// For each of the 4 segments...
for (int i = 0; i < numberOfSegments; ++i) {
    // Calculate the index where this particular segment should operate within the larger arrays.
    segmentOffset = i * segmentN;

    // Create a stream for this segment's worth of copy and work.
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    // Asynchronously copy segment's worth of pinned host memory to device over non-default stream.
    cudaMemcpyAsync(&device_array[segmentOffset], // Take care to access correct location in array.
                    &host_array[segmentOffset],   // Take care to access correct location in array.
                    segmentSize,                  // Only copy a segment's worth of memory.
                    cudaMemcpyHostToDevice,
                    stream); // Provide optional argument for non-default stream.

    // Execute segment's worth of work over same non-default stream as memory copy.
    kernel<<<number_of_blocks, threads_per_block, 0, stream>>>(&device_array[segmentOffset], segmentN);

    // `cudaStreamDestroy` will return immediately (is non-blocking), but will not actually destroy stream until
    // all stream operations are complete.
    cudaStreamDestroy(stream);
}

아래 예제 코드는 위에서 살펴본 벡터 덧셈 코드에서 벡터를 덧셈할 때 4개의 세그먼트로 분리하여 각각의 스트림에서 copy/addition을 수행하도록 변경하였습니다.

// 16_vector-add-overlap-xfer.cu
#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

__global__
void initWith(float num, float *a, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride)
    {
        a[i] = num;
    }
}

__global__
void addVectorsInto(float* result, float* a, float* b, const int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < N; i += stride) {
        result[i] = a[i] + b[i];
    }
}

void checkElementsAre(float target, float* array, const int N)
{
    for (int i = 0; i < N; i++) {
        if (array[i] != target) {
            printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
            exit(1);
        }
    }
    printf("SUCCESS! All values added correctly.\n");
}

int main()
{
    const int N = 2 << 20;
    size_t size = N * sizeof(float);

    int deviceId;
    checkCuda(cudaGetDevice(&deviceId));

    cudaDeviceProp props;
    checkCuda(cudaGetDeviceProperties(&props, deviceId));

    float *a, *b, *c, *h_c;

    checkCuda(cudaMalloc(&a, size));
    checkCuda(cudaMalloc(&b, size));
    checkCuda(cudaMalloc(&c, size));
    checkCuda(cudaMallocHost(&h_c, size));

    size_t threadsPerBlock = props.maxThreadsPerBlock;
    size_t numberOfBlocks = props.multiProcessorCount;

    cudaStream_t stream1, stream2, stream3;
    checkCuda(cudaStreamCreate(&stream1));
    checkCuda(cudaStreamCreate(&stream2));
    checkCuda(cudaStreamCreate(&stream3));

    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
    initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);

    for (int i = 0; i < 4; ++i) {
        cudaStream_t stream;
        checkCuda(cudaStreamCreate(&stream));

        addVectorsInto<<<numberOfBlocks / 4, threadsPerBlock, 0, stream>>>(&c[i * N / 4], &a[i * N / 4], &b[i * N / 4], N / 4);
        checkCuda(cudaMemcpyAsync(&h_c[i * N / 4], &c[i * N / 4], size / 4, cudaMemcpyDeviceToHost, stream));
        checkCuda(cudaStreamDestroy(stream));
    }
    checkCuda(cudaDeviceSynchronize());
    
    checkElementsAre(7, h_c, N);

    checkCuda(cudaStreamDestroy(stream1));
    checkCuda(cudaStreamDestroy(stream2));
    checkCuda(cudaStreamDestroy(stream3));

    checkCuda(cudaFree(a));
    checkCuda(cudaFree(b));
    checkCuda(cudaFree(c));
    checkCuda(cudaFreeHost(h_c));
}

위 코드를 컴파일 한 후, 프로파일링 한 결과는 다음과 같습니다.

위에서 사용된 예제 코드들 일부는 아래 github에서 확인하실 수 있습니다.

'NVIDIA > CUDA' 카테고리의 다른 글

NVIDIA Tools Extension (NVTX) (0)	2022.06.17
Multiple GPUs with CUDA C++ (0)	2022.06.16
CUDA C/C++ 기초 - (2) (0)	2022.06.13
CUDA C/C++ 기초 - (1) (1)	2022.06.10
CUDA Instructions (2) - Instruction 최적화 (0)	2022.01.28

별준

CUDA C/C++ 기초 - (3)

References

Contents

Concurrent CUDA Streams

Creating, Utilizing, and Destroying Non-Default CUDA Streams

Manual Device Memory Allocation and Copying

Manual Device Memory Management Example

Using Streams to Data Transfers

'NVIDIA > CUDA' 카테고리의 다른 글

댓글

티스토리툴바

CUDA C/C++ 기초 - (3)

References

Contents

Concurrent CUDA Streams

Creating, Utilizing, and Destroying Non-Default CUDA Streams

Manual Device Memory Allocation and Copying

Manual Device Memory Management Example

Using Streams to Data Transfers

'NVIDIA > CUDA' 카테고리의 다른 글

관련글

댓글

티스토리툴바