Làm thế nào để tự động phân bổ mảng bên trong một hạt nhân?

Question

Tôi cần phân bổ động một số mảng bên trong hàm hạt nhân. Làm thế nào tôi có thể làm điều đó?

Mã của tôi là một cái gì đó như thế:

__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. 
} 

Nhưng điều đó sẽ không làm việc. Nếu đây là bên trong mã máy chủ tôi có thể sử dụng malloc. cudaMalloc cần một con trỏ trên máy chủ và thiết bị khác trên thiết bị. Bên trong hàm hạt nhân, tôi không có con trỏ chủ.

Vì vậy, tôi nên làm gì?

Nếu mất quá lâu (một vài giây) để phân bổ tất cả các mảng (tôi cần khoảng 4 kích thước n và 5 kích thước nn), điều này sẽ không thành vấn đề. Kể từ khi hạt nhân có lẽ sẽ chạy trong 20 phút, ít nhất.

Answer 1

Phân bổ bộ nhớ động chỉ được hỗ trợ trên khả năng tính toán 2.x và phần cứng mới hơn. Bạn có thể sử dụng một trong hai C++ mới từ khóa hoặc malloc trong hạt nhân, vì vậy ví dụ của bạn có thể trở thành:

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

này phân bổ bộ nhớ trên bộ nhớ địa phương đống runtime trong đó có thời gian tồn tại của bối cảnh, do đó hãy chắc chắn rằng bạn giải phóng bộ nhớ sau khi hạt nhân kết thúc chạy nếu ý định của bạn không sử dụng lại bộ nhớ. Bạn cũng nên lưu ý rằng bộ nhớ heap thời gian chạy không thể được truy cập trực tiếp từ các API của máy chủ, vì vậy bạn không thể chuyển một con trỏ được cấp phát bên trong một hạt nhân làm đối số cho cudaMemcpy, ví dụ.

Answer 2

@talonmies đã trả lời câu hỏi của bạn về cách phân bổ động bộ nhớ trong hạt nhân. Đây là câu trả lời bổ sung, giải quyết hiệu suất của __device__ malloc() và một giải pháp thay thế mà bạn có thể muốn xem xét.

Phân bổ bộ nhớ động trong hạt nhân có thể hấp dẫn vì nó cho phép mã GPU trông giống mã CPU hơn. Nhưng nó có thể ảnh hưởng nghiêm trọng đến hiệu suất. Tôi đã viết một bài kiểm tra tự chứa và đã bao gồm nó dưới đây. Bài kiểm tra ra mắt khoảng 2,6 triệu chủ đề. Mỗi luồng sẽ điền 16 số nguyên bộ nhớ toàn cục với một số giá trị bắt nguồn từ chỉ mục chuỗi, sau đó tổng hợp các giá trị và trả về tổng.

Bài kiểm tra thực hiện hai cách tiếp cận. Cách tiếp cận đầu tiên sử dụng __device__ malloc() và cách tiếp cận thứ hai sử dụng bộ nhớ được cấp phát trước khi hạt nhân chạy.

Trên thiết bị 2.0 của tôi, hạt nhân chạy trong 1500ms khi sử dụng __device__ malloc() và 27ms khi sử dụng bộ nhớ được phân bổ trước. Nói cách khác, thử nghiệm mất 56x dài hơn để chạy khi bộ nhớ được phân bổ động trong hạt nhân. Thời gian bao gồm vòng ngoài cudaMalloc()/cudaFree(), không phải là một phần của hạt nhân. Nếu cùng một hạt nhân được khởi chạy nhiều lần với cùng một số luồng, như thường lệ, chi phí của cudaMalloc()/cudaFree() được khấu hao theo tất cả các lần khởi chạy hạt nhân. Điều đó mang lại sự khác biệt cao hơn, khoảng 60x.

Đầu cơ, tôi nghĩ rằng lần truy cập hiệu suất một phần là do việc tuần tự hóa ngầm. GPU có thể phải tuần tự hóa tất cả các cuộc gọi đồng thời tới __device__ malloc() để cung cấp các bộ nhớ riêng biệt cho từng người gọi.

Phiên bản không sử dụng __device__ malloc() phân bổ tất cả bộ nhớ GPU trước khi chạy hạt nhân. Một con trỏ tới bộ nhớ được truyền đến hạt nhân. Mỗi chuỗi tính toán chỉ mục vào bộ nhớ được cấp phát trước đó thay vì sử dụng một số __device__ malloc().

Vấn đề tiềm ẩn khi cấp phát bộ nhớ lên là, nếu chỉ một số chủ đề cần cấp phát bộ nhớ và không biết chủ đề là gì, cần phân bổ bộ nhớ cho tất cả các chuỗi. Nếu không có đủ bộ nhớ cho điều đó, nó có thể hiệu quả hơn để giảm số lượng các chủ đề cho mỗi cuộc gọi hạt nhân sau đó sử dụng __device__ malloc(). Các cách giải quyết khác có thể sẽ kết thúc việc thực hiện lại những gì __device__ malloc() đang hoạt động ở chế độ nền và sẽ có hiệu suất tương tự.

Kiểm tra việc thực hiện các __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; 
} 

Output:

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

Nếu giá trị của n và nn được biết đến trước khi kernel được gọi, thì tại sao không cudaMalloc bộ nhớ đứng về phía chủ nhà và truyền con trỏ bộ nhớ thiết bị đến hạt nhân?