सीयूडीए

Question

पर इंटर-ब्लॉक बाधा मैं सीयूडीए पर एक इंटर-ब्लॉक बाधा लागू करना चाहता हूं, लेकिन गंभीर समस्या का सामना करना चाहता हूं।

मुझे नहीं पता कि यह क्यों काम नहीं करता है।

#include <iostream> 
#include <cstdlib> 
#include <ctime> 

#define SIZE 10000000 
#define BLOCKS 100 

using namespace std; 

struct Barrier { 
    int *count; 

    __device__ void wait() { 
     atomicSub(count, 1); 
     while(*count) 
      ; 
    } 

    Barrier() { 
     int blocks = BLOCKS; 
     cudaMalloc((void**) &count, sizeof(int)); 
     cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice); 
    } 

    ~Barrier() { 
     cudaFree(count); 
    } 
}; 


__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier) 
{ 
    int tid = blockIdx.x; 

    int temp = 0; 
    while(tid < SIZE) { 
     temp += vec[tid]; 
     tid += gridDim.x; 
    } 

    cache[blockIdx.x] = temp; 

    barrier.wait(); 

    if(blockIdx.x == 0) { 
     for(int i = 0 ; i < BLOCKS; ++i) 
      *sum += cache[i]; 
    } 
} 

int main() 
{ 
    int* vec_host = (int *) malloc(SIZE * sizeof(int));  
    for(int i = 0; i < SIZE; ++i) 
     vec_host[i] = 1; 

    int *vec_dev; 
    int *sum_dev; 
    int *cache; 
    int sum_gpu = 0; 

    cudaMalloc((void**) &vec_dev, SIZE * sizeof(int)); 
    cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &sum_dev, sizeof(int)); 
    cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &cache, BLOCKS * sizeof(int)); 
    cudaMemset(cache, 0, BLOCKS * sizeof(int)); 

    Barrier barrier; 
    sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier); 

    cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost); 

    cudaFree(vec_dev); 
    cudaFree(sum_dev); 
    cudaFree(cache); 
    free(vec_host); 
    return 0; 
} 

वास्तव में, यहां तक ​​कि अगर मैं इंतजार पुनर्लेखन() निम्नलिखित

__device__ void wait() { 
     while(*count != 234124) 
      ; 
    } 

कार्यक्रम के रूप में सामान्य रूप से बाहर निकालता है। लेकिन मुझे इस मामले में एक अनंत लूप मिलने की उम्मीद है।

Answer 1

दुर्भाग्यवश, आप जो हासिल करना चाहते हैं (इंटर-ब्लॉक संचार/सिंक्रनाइज़ेशन) CUDA में सख्ती से संभव नहीं है। सीयूडीए प्रोग्रामिंग गाइड कहता है कि "थ्रेड ब्लॉक को स्वतंत्र रूप से निष्पादित करने की आवश्यकता होती है: समानांतर या श्रृंखला में किसी भी क्रम में उन्हें निष्पादित करना संभव होना चाहिए।" इस प्रतिबंध का कारण थ्रेड ब्लॉक शेड्यूलर में लचीलापन की अनुमति देना है, और कोड को संख्याओं के साथ अज्ञात रूप से स्केल करने की अनुमति देना है। एकमात्र समर्थित इंटर-ब्लॉक सिंक्रनाइज़ेशन विधि एक और कर्नेल लॉन्च करना है: कर्नेल लॉन्च (उसी स्ट्रीम के भीतर) निहित सिंक्रनाइज़ेशन पॉइंट्स हैं।

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

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

अपने आवेदन के लिए, एक समाधान खोजने का प्रयास करें जो अंतर-ब्लॉक सिंक्रनाइज़ेशन पर निर्भर नहीं है, क्योंकि (CUDA प्रोग्रामिंग मॉडल में एक संकेत परिवर्तन को छोड़कर) यह संभव नहीं है।

Answer 2

संकलक अनुकूलन समस्या की तरह दिखता है। मैं PTX-कोड को पढ़ने के साथ अच्छा नहीं हूँ, लेकिन ऐसा लगता है संकलक बिल्कुल while -loop नहीं दिखाए हैं (यहां तक ​​कि जब -O0 साथ संकलित) दिखता है:

.loc 3 41 0 
cvt.u64.u32  %rd7, %ctaid.x; // Save blockIdx.x to rd7 
ld.param.u64 %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache]; 
mov.s32  %r8, %ctaid.x; // Now calculate ouput address 
mul.wide.u32 %rd9, %r8, 4; 
add.u64  %rd10, %rd8, %rd9; 
st.global.s32 [%rd10+0], %r5; // Store result to cache[blockIdx.x] 
.loc 17 128 0 
ld.param.u64 %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11 
mov.s32  %r9, -1; // put -1 to r9 
atom.global.add.s32  %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused) 
cvt.u32.u64  %r11, %rd7; // Put blockIdx.x saved in rd7 to r11 
mov.u32  %r12, 0; // Put 0 to r12 
setp.ne.u32  %p3, %r11, %r12; // if(blockIdx.x == 0) 
@%p3 bra $Lt_0_5122; 
ld.param.u64 %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum]; 
ld.global.s32 %r13, [%rd12+0]; 
mov.s64  %rd13, %rd8; 
mov.s32  %r14, 0; 

सीपीयू कोड के मामले में, इस तरह के व्यवहार को रोका जाता है परिवर्तनीय को volatile उपसर्ग के साथ घोषित करके। लेकिन फिर भी अगर हम घोषित count रूप int __device__ count (और उचित रूप से कोड बदलने), volatile विनिर्देशक जोड़ने सिर्फ संकलन टूट जाता है (त्रुटियों लोके argument of type "volatile int *" is incompatible with parameter of type "void *" के साथ)

मैं CUDA एसडीके से threadFenceReduction उदाहरण को सुझाव देते हैं। वहां वे आपके जैसा ही कर रहे हैं, लेकिन अंतिम सारांश करने के लिए ब्लॉक को पूर्वनिर्धारित के बजाए रनटाइम में चुना जाता है, और while -loop समाप्त हो गया है, क्योंकि वैश्विक चर पर स्पिन-लॉक बहुत धीमा होना चाहिए।

Answer 3

सिंक्रनाइज़ेशन को अवरुद्ध करने के लिए ब्लॉक संभव है। यह paper देखें।
पेपर यह कैसे काम करता है इस बारे में बहुत विस्तार से नहीं जाता है, लेकिन यह __syncthreads() के संचालन पर निर्भर करता है; वर्तमान ब्लॉक के लिए विराम-बाधा बनाने के लिए, ... अन्य ब्लॉक को सिंक बिंदु पर जाने के लिए प्रतीक्षा करते समय।

पेपर में नोट नहीं किया गया एक आइटम यह है कि सिंक केवल तभी संभव है जब ब्लॉक की संख्या पर्याप्त छोटी हो या एसएम की संख्या हाथ पर काम के लिए पर्याप्त हो। यानी यदि आपके पास 4 एसएम हैं और 5 ब्लॉक सिंक करने की कोशिश कर रहे हैं, तो कर्नेल डेडलॉक होगा।

उनके दृष्टिकोण के साथ, मैं कई ब्लॉक के बीच एक लंबा सीरियल कार्य फैलाने में सक्षम हूं, आसानी से एक ब्लॉक दृष्टिकोण पर 30% समय बचा रहा हूं। यानी ब्लॉक-सिंक मेरे लिए काम करता है।