opencl
कर्नेल मूल बातें

परिचय

ग्रेस्केल कर्नेल

ग्रेस्केल छवि बनाने के लिए कर्नेल का निर्माण करें। हम छवि डेटा का उपयोग करेंगे जो कि प्रत्येक घटक के लिए यूएनजीएस और ऑर्डर RGBA के साथ परिभाषित किया गया है।

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

अब उस कोड स्टेप बाई स्टेप चलते हैं। पहली पंक्ति प्रकार के नमूने के __constant मेमोरी क्षेत्र में एक चर बनाती है। इस नमूने का उपयोग हमारी छवि डेटा तक पहुंच को निर्दिष्ट करने के लिए किया जाता है। कृपया एक पूर्ण प्रलेखन के लिए क्रोनोस डॉक्स देखें।

हमने इनपुट को read_only के रूप में और आबंटित के रूप में आबंटित किया है इससे पहले कि हम अपने कर्नेल को कहते हैं, इसलिए हम उन संशोधक को यहां जोड़ते हैं।

image2d और image3d को हमेशा वैश्विक मेमोरी पर आवंटित किया जाता है, इसलिए हम यहां __global संशोधक को छोड़ सकते हैं।

फिर हम अपनी थ्रेड आईडी प्राप्त करते हैं जो पिक्सेल को परिभाषित करता है जिसे हम ग्रेस्केल में बदलने जा रहे हैं। हम यह भी सुनिश्चित करने के लिए आकार को क्वेरी करते हैं कि हमारा थ्रेड असंबद्ध मेमोरी तक नहीं पहुंच रहा है। यदि आप भूल जाते हैं तो यह आपके कर्नेल को क्रैश कर देगा।

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

कर्नेल स्केलेटन

चलो सबसे सरल गिरी के माध्यम से चलते हैं और इसके कुछ रूपांतर हैं

__kernel void myKernel() {
}

एक कर्नेल जिसे आपके मुख्य कोड से शुरू किया जा सकता है, उसे __kernel कीवर्ड द्वारा पहचाना जाता है। कर्नेल फ़ंक्शन में केवल वापसी प्रकार शून्य हो सकता है।

__kernel void myKernel(float a, uint b, byte c) {

}

बेशक आप अधिक कार्य बना सकते हैं जो गुठली के रूप में उजागर नहीं होते हैं। इस स्थिति में आप केवल __kernel संशोधक को छोड़ सकते हैं।

एक फ़ंक्शन किसी अन्य C / C ++ फ़ंक्शन की तरह चर को उजागर कर सकता है। फर्क सिर्फ इतना है कि जब आप मेमोरी को रेफर करना चाहते हैं। यह सभी बिंदुओं पर लागू होता है कि क्या वे तर्क हैं या कोड में उपयोग किए जाते हैं।

float*  ptr;

एक स्मृति क्षेत्र के लिए एक संकेतक है केवल निष्पादन धागा तक पहुंच है। वास्तव में यह वैसा ही है

__private float* ptr;

चार अलग-अलग मेमोरी क्षेत्र के मॉडिफायर उपलब्ध हैं। कर्नेल के अंदर आपको आमतौर पर इसके बारे में चिंता करने की ज़रूरत नहीं है, लेकिन तर्कों के लिए यह आवश्यक है।

• __global: यह संशोधक एक पॉइंटर को संदर्भित करता है जिसे वैश्विक मेमोरी में रखा गया है
• __constant: एक निरंतर मेमोरी पॉइंटर को संदर्भित करता है
• __local: एक साझा मेमोरी पॉइंटर को संदर्भित करता है
• __पेयर: एक स्थानीय मेमोरी पॉइंटर को संदर्भित करता है

इसके अलावा हम परिभाषित कर सकते हैं कि हम मेमोरी को कैसे एक्सेस करना चाहते हैं

• कोई संशोधक नहीं: पढ़ें और लिखें
• __सिफ़ पढ़िये
• __write_only

उन झंडों का मिलान उस तरह से करना है जैसा हमने अपने मेजबान पर मेमोरी बफर को वापस आवंटित किया था।

कर्नेल आईडी

डेटा के साथ ठीक से काम करने के लिए प्रत्येक थ्रेडब्लॉक / ग्लोबल थ्रेड पूल में अपनी स्थिति जानना आवश्यक है। इसके साथ संग्रह किया जा सकता है

get_local_id($dim);
get_global_id($dim);

वे दो कार्य थ्रेडब्लॉक या सभी थ्रेड्स के सापेक्ष थ्रेड की स्थिति लौटाते हैं।

get_working_dim();

कर्नेल के साथ लॉन्च किए गए आयामों की कुल संख्या हो जाती है।

get_local_size($dim);
get_global_size($dim);

थ्रेडब्लॉक में या दिए गए आयाम के लिए कुल थ्रेड की कुल संख्या प्राप्त होती है।

कैविएट: हमेशा सुनिश्चित करें कि आपका थ्रेड आपके डेटा आकार से अधिक नहीं है। ऐसा होने की बहुत संभावना है और इसके लिए हमेशा जाँच की जानी चाहिए।

ओपनसीएल में वैक्टर

प्रत्येक मौलिक opencl प्रकार का एक वेक्टर संस्करण है। आप प्रकार के बाद वांछित घटकों की संख्या जोड़कर वेक्टर प्रकार का उपयोग कर सकते हैं। घटकों की समर्थित संख्या 2,3,4,8 और 16 है। ओपनसीएल 1.0 तीन घटकों की पेशकश नहीं करता है।

आप दो तरीकों का उपयोग करके किसी भी वेक्टर को आरंभीकृत कर सकते हैं:

• एक एकल स्केलर प्रदान करें
• सभी घटकों को संतुष्ट करें

float4 a = (float4)(1); //a = (1, 1, 1, 1)

या

float4 b = (float4)(1, 2, 3, 4);
float4 c = (float4)(1, (float3)(2));

या वैक्टर के किसी भी अन्य संयोजन जो घटकों की संख्या को संतुष्ट करते हैं। एक वेक्टर के तत्वों तक पहुंचने के लिए आप विभिन्न तरीकों का उपयोग कर सकते हैं। आप या तो अनुक्रमण का उपयोग कर सकते हैं:

a[0] = 2;

या शाब्दिक उपयोग करें। शाब्दिक का लाभ यह है कि आप उन्हें किसी भी तरह से जोड़ सकते हैं, जो आप चाहते हैं, एक पल में। आप सभी वेक्टर घटकों को एक्सेस कर सकते हैं

a.s0 = 2; // same as a[0] = 2

आप एक नए वेक्टर में कई घटकों को जोड़ सकते हैं

a.s02 = (float2)(0, 0); // same as  a[0] = 0; a[2] = 0; or even a.s20 = (float2)(0, 0)

आप किसी भी तरह से घटकों के क्रम को बदल सकते हैं।

a.s1423 = a.s4132; // flip the vector

लेकिन आप ऐसा कुछ नहीं कर सकते

a.s11 = ... // twice the same component is not possible

वेक्टर घटकों तक पहुंचने के लिए कुछ सुविधाजनक आशुलिपि हैं। निम्नलिखित आशुलिपि केवल 2, 4, 8 और 16 के आकारों पर लागू होती हैं

a.hi //=a.s23 for vectors of size 4, a.4567 for size 8 and so on.
a.lo //=a.s01
a.even //=a.s02
a.odd //=a.13

वेक्टर आकार 2,3 और 4 के लिए कुछ अतिरिक्त शॉर्टहैंड हैं

a.x //=a.s0
a.y //=a.s1
a.z //=a.s2
a.w //=a.s3

गामा सुधार कर्नेल

चलो एक गामा सुधार कर्नेल को देखते हैं

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Gamma(__read_only image2d_t input, __write_only image2d_t output, __constant float gamma) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color = pow(color, (float4)(gamma));
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

अब उस कोड स्टेप बाई स्टेप चलते हैं। पहली पंक्ति प्रकार के नमूने के __constant मेमोरी क्षेत्र में एक चर बनाती है। इस नमूने का उपयोग हमारी छवि डेटा तक पहुंच को निर्दिष्ट करने के लिए किया जाता है। कृपया एक पूर्ण प्रलेखन के लिए क्रोनोस डॉक्स देखें।

हमने इनपुट को read_only के रूप में और आबंटित के रूप में आबंटित किया है इससे पहले कि हम अपने कर्नेल को कहते हैं, इसलिए हम उन संशोधक को यहां जोड़ते हैं।

image2d और image3d को हमेशा वैश्विक मेमोरी पर आवंटित किया जाता है, इसलिए हम यहां __global संशोधक को छोड़ सकते हैं। हमारा गामा मूल्य __constant मेमोरी में स्थित है, इसलिए हम इसे भी निर्दिष्ट करते हैं।

फिर हम अपनी थ्रेड आईडी प्राप्त करते हैं जो पिक्सेल को परिभाषित करता है जिसे हम सही तरीके से गामा करने जा रहे हैं। हम यह भी सुनिश्चित करने के लिए आकार को क्वेरी करते हैं कि हमारा थ्रेड असंबद्ध मेमोरी तक नहीं पहुंच रहा है। यदि आप भूल जाते हैं तो यह आपके कर्नेल को क्रैश कर देगा।

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