커널 내부에 동적으로 배열을 할당하는 방법은 무엇입니까?

Question

커널 함수 내에 일부 배열을 동적으로 할당해야합니다. 내가 어떻게 그럴 수 있니?

내 코드는 같은 것입니다 :

__global__ func(float *grid_d,int n, int nn){ 
    int i,j; 
    float x[n],y[nn]; 
    //Do some really cool and heavy computations here that takes hours. 
} 

그러나 그것은 작동하지 않습니다. 이것이 호스트 코드 안에 있다면 malloc을 사용할 수 있습니다. cudaMalloc은 호스트에 포인터가 필요하고 다른 장치에는 포인터가 필요합니다. 커널 함수 안에는 호스트 포인터가 없습니다.

그럼 어떻게해야합니까?

만약 모든 어레이를 할당하는 데 너무 오래 걸리는 경우 (크기 n과 크기 nn 중 약 4 개가 필요함)이 문제는 발생하지 않습니다. 커널은 아마도 적어도 20 분 동안 작동 할 것입니다.

Answer 1

동적 메모리 할당은 컴퓨팅 기능 2.x 이상에서만 지원됩니다. 당신은 당신의 예제가 될 수 있도록 커널에서 C++ 새로운 키워드 또는 malloc에 ​​하나를 사용할 수 있습니다

이

__global__ func(float *grid_d,int n, int nn){ 
    int i,j; 
    float *x = new float[n], *y = new float[nn]; 
} 

이 문맥의 수명을 가지고 로컬 메모리 런타임 힙에 메모리를 할당, 그래서 당신이 무료 확인 메모리를 다시 사용하지 않으려는 경우 커널이 실행을 마친 후 메모리. 또한 런타임 힙 메모리는 호스트 API에서 직접 액세스 할 수 없으므로 예를 들어 커널 내부에 할당 된 포인터를 인수로 cudaMemcpy에 전달할 수 없습니다.

Answer 2

@talonmies가 커널 내에서 동적으로 메모리를 할당하는 방법에 대한 질문에 답변했습니다. 이는 보충 답변으로, __device__ malloc()의 성능 및 고려해야 할 대안에 대해 설명합니다.

커널에서 동적으로 메모리를 할당하면 GPU 코드가 CPU 코드와 비슷하게 보이기 때문에 유혹을받을 수 있습니다. 그러나 성능에 심각한 영향을 줄 수 있습니다. 필자는 자체 시험을 작성하여 아래에 포함 시켰습니다. 이 테스트는 약 260 만 개의 스레드를 시작합니다. 각 스레드는 전역 메모리의 16 정수를 스레드 인덱스에서 파생 된 일부 값으로 채운 다음 값을 합하여 합계를 반환합니다.

이 테스트에서는 두 가지 접근 방식을 구현합니다. 첫 번째 방법은 __device__ malloc()을 사용하고 두 번째 방법은 커널을 실행하기 전에 할당 된 메모리를 사용합니다.

내 2.0 장치에서 커널은 미리 할당 된 메모리를 사용할 때 __device__ malloc() 및 27ms를 사용할 때 1500ms에서 실행됩니다. 즉, 메모리가 커널 내에서 동적으로 할당 될 때 이 56x 길어질 때이 실행됩니다. 이 시간에는 커널의 일부가 아닌 외부 루프 cudaMalloc()/cudaFree()이 포함됩니다. 종종 같은 수의 동일한 스레드로 같은 커널을 여러 번 실행하면 cudaMalloc()/cudaFree()의 비용이 모든 커널 시작에 대해 상각됩니다. 그 차이는 약 60 배까지 훨씬 커집니다.

추측 해 보면 성능 ​​저하는 부분적으로 암시 적 직렬화에 의한 것이라고 생각합니다. GPU는 각 호출자에게 별도의 메모리 청크를 제공하기 위해 모든 동시 호출을 __device__ malloc()으로 직렬화해야합니다.

__device__ malloc()을 사용하지 않는 버전은 커널을 실행하기 전에 모든 GPU 메모리를 할당합니다. 메모리에 대한 포인터가 커널에 전달됩니다. 각 스레드는 __device__ malloc()을 사용하는 대신 이전에 할당 된 메모리에 대한 인덱스를 계산합니다.

메모리를 할당 할 때 잠재적 인 문제는 일부 스레드 만 메모리를 할당해야하며 스레드가 어떤 스레드인지 알 수없는 경우 모든 스레드에 메모리를 할당해야한다는 것입니다. 충분한 메모리가 없으면 커널 호출 당 스레드 수를 줄이고 __device__ malloc()을 사용하는 것이 더 효율적일 수 있습니다. 다른 해결 방법은 아마도 __device__ malloc()이 백그라운드에서 수행중인 작업을 다시 구현하는 결과를 가져오고 비슷한 성능 저하를 보게됩니다.

테스트 __device__ malloc()의 성능 :

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <stdio.h> 

const int N_ITEMS(16); 

#define USE_DYNAMIC_MALLOC 

__global__ void test_malloc(int* totals) 
{ 
    int tx(blockIdx.x * blockDim.x + threadIdx.x); 

    int* s(new int[N_ITEMS]); 

    for (int i(0); i < N_ITEMS; ++i) { 
    s[i] = tx * i; 
    } 

    int total(0); 
    for (int i(0); i < N_ITEMS; ++i) { 
    total += s[i]; 
    } 

    totals[tx] = total; 

    delete[] s; 
} 

__global__ void test_malloc_2(int* items, int* totals) 
{ 
    int tx(blockIdx.x * blockDim.x + threadIdx.x); 

    int* s(items + tx * N_ITEMS); 

    for (int i(0); i < N_ITEMS; ++i) { 
    s[i] = tx * i; 
    } 

    int total(0); 
    for (int i(0); i < N_ITEMS; ++i) { 
    total += s[i]; 
    } 

    totals[tx] = total; 
} 

int main() 
{ 
    cudaError_t cuda_status; 

    cudaSetDevice(0); 

    int blocks_per_launch(1024 * 10); 
    int threads_per_block(256); 

    int threads_per_launch(blocks_per_launch * threads_per_block); 

    int* totals_d; 
    cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int)); 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    cudaDeviceSynchronize(); 
    cudaEventRecord(start, 0); 

#ifdef USE_DYNAMIC_MALLOC 
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int)); 

    test_malloc<<<blocks_per_launch, threads_per_block>>>(totals_d); 
#else 
    int* items_d; 
    cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS); 

    test_malloc_2<<<blocks_per_launch, threads_per_block>>>(items_d, totals_d); 

    cudaFree(items_d); 
#endif 

    cuda_status = cudaDeviceSynchronize(); 
    if (cuda_status != cudaSuccess) { 
    printf("Error: %d\n", cuda_status); 
    exit(1); 
    } 

    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime, start, stop); 

    printf("Elapsed: %f\n", elapsedTime); 

    int* totals_h(new int[threads_per_launch]); 
    cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost); 
    if (cuda_status != cudaSuccess) { 
    printf("Error: %d\n", cuda_status); 
    exit(1); 
    } 

    for (int i(0); i < 10; ++i) { 
    printf("%d ", totals_h[i]); 
    } 
    printf("\n"); 

    cudaFree(totals_d); 
    delete[] totals_h; 

    return cuda_status; 
} 

출력 : N의 값과 윈이 알려진다면

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe 
Elapsed: 27.311169 
0 120 240 360 480 600 720 840 960 1080 

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe 
Elapsed: 1516.711914 
0 120 240 360 480 600 720 840 960 1080 

Answer 3

커널을 호출하기 전에, 왜 호스트 측의 메모리를 cudaMalloc하지 그리고 장치 메모리 포인터를 커널에 건네 줍니까?