गतिशील रूप से कर्नेल के अंदर सरणी आवंटित करने के लिए कैसे?

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 को होस्ट पर एक पॉइंटर और डिवाइस पर अन्य की आवश्यकता होती है। कर्नेल फ़ंक्शन के अंदर मेरे पास होस्ट पॉइंटर नहीं है।

तो, मुझे क्या करना चाहिए?

यदि सभी सरणी आवंटित करने में बहुत लंबा (कुछ सेकंड) लगता है (मुझे आकार एन के आकार 4 और आकार के 5 की आवश्यकता है), यह कोई समस्या नहीं होगी। चूंकि कर्नेल शायद कम से कम 20 मिनट तक चलाएगा।

Answer 1

गतिशील स्मृति आवंटन केवल गणना क्षमता 2.x और नए हार्डवेयर पर समर्थित है। तुम्हें पता है, कर्नेल में या तो सी ++ नया कीवर्ड या malloc उपयोग करती हैं इसलिए अपने उदाहरण बन सकता है कर सकते हैं:

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

यह एक स्थानीय स्मृति क्रम ढेर जो संदर्भ के जीवनकाल है पर स्मृति आबंटित करता है, ताकि आप मुक्त कर यह सुनिश्चित कर लें कर्नेल के चलने के बाद स्मृति अगर आपका इरादा फिर से मेमोरी का उपयोग नहीं करना है। आपको यह भी ध्यान रखना चाहिए कि रनटाइम हीप मेमोरी को मेजबान एपीआई से सीधे एक्सेस नहीं किया जा सकता है, इसलिए आप कर्नेल के अंदर आवंटित पॉइंटर को cudaMemcpy पर तर्क के रूप में पास नहीं कर सकते हैं।

Answer 2

@talonmies ने कर्नेल के भीतर गतिशील रूप से स्मृति आवंटित करने के तरीके पर आपके प्रश्न का उत्तर दिया। यह एक पूरक उत्तर के रूप में है, __device__ malloc() के प्रदर्शन को संबोधित करना और एक विकल्प जिसे आप विचार करना चाहते हैं।

कर्नेल में गतिशील रूप से स्मृति आवंटित करना आकर्षक हो सकता है क्योंकि यह GPU कोड को CPU कोड की तरह दिखने की अनुमति देता है। लेकिन यह प्रदर्शन को गंभीरता से प्रभावित कर सकता है। मैंने स्वयं निहित परीक्षण लिखा और इसे नीचे शामिल किया है। परीक्षण में कुछ 2.6 मिलियन धागे लॉन्च किए गए। प्रत्येक थ्रेड थ्रेड इंडेक्स से प्राप्त कुछ मानों के साथ वैश्विक मेमोरी के 16 पूर्णांक को पॉप्युलेट करता है, फिर मानों को जोड़ता है और योग देता है।

परीक्षण दो दृष्टिकोण लागू करता है। पहला दृष्टिकोण __device__ malloc() का उपयोग करता है और दूसरा दृष्टिकोण कर्नेल चलाने से पहले आवंटित स्मृति का उपयोग करता है।

मेरे 2.0 डिवाइस पर, पूर्व-आवंटित स्मृति का उपयोग करते समय __device__ malloc() और 27ms का उपयोग करते समय कर्नेल 1500ms में चलता है। दूसरे शब्दों में, परीक्षण को 56x लंबा चलाता है जब स्मृति को कर्नेल के भीतर गतिशील रूप से आवंटित किया जाता है। उस समय बाहरी लूप cudaMalloc()/cudaFree() शामिल है, जो कर्नेल का हिस्सा नहीं है। एक ही गिरी, धागे की एक ही नंबर के साथ कई बार शुरू किया गया है, तो जैसा कि अक्सर मामला है, cudaMalloc()/cudaFree() की लागत सभी गिरी की शुरूआत से अधिक परिशोधित कर रहा है। इससे अंतर लगभग 60x तक पहुंच जाता है।

अनुमान लगाते हुए, मुझे लगता है कि प्रदर्शन हिट अंतर्निहित धारावाहिकता के कारण है। जीपीयू को प्रत्येक कॉलर को स्मृति के अलग-अलग हिस्सों को प्रदान करने के लिए __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; 
} 

आउटपुट:

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

हैं nn n के मूल्य और जाने जाते थे गिरी कहा जाता है से पहले, तो क्यों मेजबान पक्ष पर स्मृति cudaMalloc नहीं और कर्नेल को डिवाइस मेमोरी पॉइंटर में पास करें?