Double buffering

#include <cuda_runtime.h>
#include <iostream>

#define DATA_SIZE 1024 * 1024  // 데이터 크기
#define CHUNK_SIZE 256 * 1024 // 버퍼 크기

__global__ void processKernel(float *data, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        data[idx] = data[idx] * 2.0f; // 간단한 계산
    }
}

void doubleBufferingExample() {
    // 전체 데이터를 할당
    float *hostData = new float[DATA_SIZE];
    for (int i = 0; i < DATA_SIZE; i++) {
        hostData[i] = i * 1.0f;
    }

    // GPU 메모리 버퍼 2개 생성
    float *devBuffer1, *devBuffer2;
    cudaMalloc(&devBuffer1, CHUNK_SIZE * sizeof(float));
    cudaMalloc(&devBuffer2, CHUNK_SIZE * sizeof(float));

    // 스트림 2개 생성
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);

    // Double Buffering
    int chunks = DATA_SIZE / CHUNK_SIZE;
    for (int i = 0; i < chunks; i++) {
        // 현재와 다음 스트림 선택
        cudaStream_t currentStream = (i % 2 == 0) ? stream1 : stream2;
        float *currentBuffer = (i % 2 == 0) ? devBuffer1 : devBuffer2;

        // 데이터 복사: CPU -> GPU
        cudaMemcpyAsync(currentBuffer, &hostData[i * CHUNK_SIZE], CHUNK_SIZE * sizeof(float),
                        cudaMemcpyHostToDevice, currentStream);

        // 커널 실행
        processKernel<<<(CHUNK_SIZE + 255) / 256, 256, 0, currentStream>>>(currentBuffer, CHUNK_SIZE);

        // 데이터 복사: GPU -> CPU
        cudaMemcpyAsync(&hostData[i * CHUNK_SIZE], currentBuffer, CHUNK_SIZE * sizeof(float),
                        cudaMemcpyDeviceToHost, currentStream);
    }

    // 스트림 동기화
    cudaStreamSynchronize(stream1);
    cudaStreamSynchronize(stream2);

    // 결과 확인
    for (int i = 0; i < 10; i++) {
        std::cout << hostData[i] << " ";
    }
    std::cout << std::endl;

    // 메모리 해제
    delete[] hostData;
    cudaFree(devBuffer1);
    cudaFree(devBuffer2);
    cudaStreamDestroy(stream1);
    cudaStreamDestroy(stream2);
}

int main() {
    doubleBufferingExample();
    return 0;
}