병렬기법
홈 > CUDA > 병렬기법 > OpenMP+CUDA

OpenMP+CUDA

 OpenMP 와 CUDA를 동시에 사용하여 MultiGPU를 제어할 수 있습니다.

 

다음은 NVIDIA에서 제공하는 CUDA SDK의 예제입니다. GPU 갯수를 인식하여 OpenMP 병렬구문을 사용합니다.

 

 

새로운 프로젝트에 OpenMP + CUDA를 사용하는 것에 대한 문의가 급증하고 있어서 추가 설정 내용 언급해 드리겠습니다. CUDA SDK예제는 OpenMP 세팅이 미리 정의되어있습니다. 새로 작성시는 프로젝트 속성에 추가적인 OpenMP 설정을 하셔야 합니다. -openmp /openmp 등의 명령어를 추가하셔야 하는데, 리눅스, 윈도우즈에 따라 다릅니다. 좀더 자세한 내용은 자료를 올려드리도록 하겠습니다.

 

#include

#include     

#include

 

using namespace std;

 

__global__ void kernelAddConstant(int *g_a, const int b)

{

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    g_a[idx] += b;

}

 

 

int correctResult(int *data, const int n, const int b)

{

        for(int i = 0; i < n; i++)

               if(data[i] != i + b)

                       return 0;

        return 1;

}

 

int main(int argc, char *argv[])

{

        int num_gpus = 0;     

 

    cudaGetDeviceCount(&num_gpus);

        if(num_gpus < 1)

        {

               printf("no CUDA capable devices were detected\n");

               return 1;

        }

 

    printf("number of host CPUs:\t%d\n", omp_get_num_procs());

    printf("number of CUDA devices:\t%d\n", num_gpus);

    for(int i = 0; i < num_gpus; i++)

    {

        cudaDeviceProp dprop;

        cudaGetDeviceProperties(&dprop, i);

               printf("   %d: %s\n", i, dprop.name);

    }

        printf("---------------------------\n");

 

unsigned int n = num_gpus * 8192;

    unsigned int nbytes = n * sizeof(int);

        int *a = 0;            int b = 3;            

        a = (int*)malloc(nbytes);

        if(0 == a)

        {

               printf("couldn't allocate CPU memory\n");

               return 1;

        }

        for(unsigned int i = 0; i < n; i++)

        a[i] = i;

   

        omp_set_num_threads(num_gpus);

#pragma omp parallel

    {         

        unsigned int cpu_thread_id = omp_get_thread_num();

               unsigned int num_cpu_threads = omp_get_num_threads();

               int gpu_id = -1;

               CUDA_SAFE_CALL(cudaSetDevice(cpu_thread_id % num_gpus));                             CUDA_SAFE_CALL(cudaGetDevice(&gpu_id));

               printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id, num_cpu_threads, gpu_id);

 

               int *d_a = 0; 

               int *sub_a = a + cpu_thread_id * n / num_cpu_threads;       

               unsigned int nbytes_per_kernel = nbytes / num_cpu_threads;

               dim3 gpu_threads(128);

               dim3 gpu_blocks(n / (gpu_threads.x * num_cpu_threads));

 

        CUDA_SAFE_CALL(cudaMalloc((void**)&d_a, nbytes_per_kernel));

        CUDA_SAFE_CALL(cudaMemset(d_a, 0, nbytes_per_kernel));

        CUDA_SAFE_CALL(cudaMemcpy(d_a, sub_a, nbytes_per_kernel, cudaMemcpyHostToDevice));

        kernelAddConstant<<>>(d_a, b);

        CUDA_SAFE_CALL(cudaMemcpy(sub_a, d_a, nbytes_per_kernel, cudaMemcpyDeviceToHost));

 

        CUDA_SAFE_CALL(cudaFree(d_a));

 

               CUDA_SAFE_CALL(cudaThreadExit());

    }

        printf("---------------------------\n");

       

        if(cudaSuccess != cudaGetLastError())

               printf("%s\n", cudaGetErrorString(cudaGetLastError()));

 

    if(correctResult(a, n, b))

        printf("Test PASSED\n");

    else

        printf("Test FAILED\n");

 

    free(a);  

 

        CUT_EXIT(argc, argv);

 

    return 0;

}