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

परिचय

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

मैं एक छोटा सा कोड दिखा कर शुरू करूंगा और समझाऊंगा कि वहां क्या हो रहा है, जबकि OpenCL कैसे काम करता है, इस बारे में बात कर रहा हूं।

सबसे पहले, मैं एक बल्कि तुच्छ कोड का वर्णन करूंगा और जो लोग OpenCL के ~~जादू~~ को देखने के लिए उत्सुक नहीं हैं, वे पहले भाग को छोड़ सकते हैं (बस अंतिम पैराग्राफ पढ़ें जहां मैं MathCalculations फ़ंक्शन का वर्णन करता हूं, यह महत्वपूर्ण है। और यदि आप OpenCL के बारे में जानते हैं और आप परीक्षा परिणाम देखना चाहते हैं, तो सीधे पांचवें खंड पर जाएं, लेकिन वैसे भी MathCalculations को देखें)।

int main (int argc, char * argv [])

int main(int argc, char* argv[]) { GenerateTestData(); PerformCalculationsOnHost(); //Get all available platforms vector<cl::Platform> platforms; cl::Platform::get(&platforms); for (int iPlatform=0; iPlatform<platforms.size(); iPlatform++) { //Get all available devices on selected platform std::vector<cl::Device> devices; platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices); //Perform test on each device for (int iDevice=0; iDevice<devices.size(); iDevice++) { try { PerformTestOnDevice(devices[iDevice]); } catch(cl::Error error) { std::cout << error.what() << "(" << error.err() << ")" << std::endl; } CheckResults(); } } //Clean buffers delete[](pInputVector1); delete[](pInputVector2); delete[](pOutputVector); delete[](pOutputVectorHost); return 0; }

यह है कि OpenCL के परीक्षण के लिए मेरे छोटे कार्यक्रम का मुख्य रूप, या बल्कि, एक निश्चित अमूर्त गणितीय अभिव्यक्ति की गणना करना, जो हम बाद में आएंगे। तो, चलो लाइन को समझने के लिए कि यहां क्या हो रहा है।

भाग एक - स्रोत डेटा और कम्प्यूटिंग की पारंपरिक विधि की शुरुआत

GenerateTestData (); यह असाधारण कुछ भी नहीं करता है, लेकिन बस इनपुट और आउटपुट सरणियों के लिए मेमोरी आवंटित करता है, और यादृच्छिक डेटा के साथ इनपुट सरणियों को भी भरता है।

शून्य GenerateTestData ()

 void GenerateTestData() { pInputVector1 = new float[DATA_SIZE]; pInputVector2 = new float[DATA_SIZE]; pOutputVector = new float[DATA_SIZE]; pOutputVectorHost = new float[DATA_SIZE]; srand (time(NULL)); for (int i=0; i<DATA_SIZE; i++) { pInputVector1[i] = rand() * 1000.0 / RAND_MAX; pInputVector2[i] = rand() * 1000.0 / RAND_MAX; } }

अगला थोड़ा और दिलचस्प समारोह आता है:

शून्य प्रदर्शन

 void PerformCalculationsOnHost() { cout << "Device: Host" << endl << endl; //Some performance measurement timeValues.clear(); __int64 start_count; __int64 end_count; __int64 freq; QueryPerformanceFrequency((LARGE_INTEGER*)&freq); for(int iTest=0; iTest<(TESTS_NUMBER/10); iTest++) { QueryPerformanceCounter((LARGE_INTEGER*)&start_count); for(int iJob=0; iJob<DATA_SIZE; iJob++) { //Check boundary conditions if (iJob >= DATA_SIZE) break; //Perform calculations pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]); } QueryPerformanceCounter((LARGE_INTEGER*)&end_count); double time = 1000 * (double)(end_count - start_count) / (double)freq; timeValues.push_back(time); } hostPerformanceTimeMS = std::accumulate(timeValues.begin(), timeValues.end(), 0)/timeValues.size(); PrintTimeStatistic(); }

इसमें पहला चक्र

 for(int iTest=0; iTest<(TESTS_NUMBER/10); iTest++)

अधिक सटीक रनटाइम प्राप्त करने के लिए कई बार परीक्षण चलाने की आवश्यकता होती है। प्रत्येक परीक्षण की गणना का समय टाइमवैल्यूज़ एरे में संग्रहीत किया जाता है, जहाँ से औसत मूल्य की गणना की जाती है और hostPerformingTimeMS में संग्रहीत की जाती है।

दूसरा चक्र

 for(int iJob=0; iJob<DATA_SIZE; iJob++)

क्रमिक रूप से इनपुट सरणियों के तत्वों पर कुछ गणितीय गणना करता है और उन्हें आउटपुट सरणी में संग्रहीत करता है।

जैसा कि हम देख सकते हैं, इस कोड में कुछ भी असामान्य नहीं है; यह एक नियमित सिस्टम कंपाइलर द्वारा संकलित किया जाता है और केंद्रीय प्रोसेसर पर क्रमिक रूप से निष्पादित होता है, अधिकांश कोड की तरह जो हम सभी हर दिन लिखते हैं। और हमें उसे ओपनसीएल द्वारा प्राप्त परिणामों की तुलना करने के लिए उसकी आवश्यकता है, और यह भी समझने के लिए कि हमें किस प्रकार का प्रदर्शन प्राप्त होता है।

तुरंत आपको MathCalculations में देखना चाहिए और देखना चाहिए कि सब कुछ पूरी तरह से उबाऊ है:

फ्लोट MathCalculations (फ्लोट ए, फ्लोट बी)

 float MathCalculations(float a, float b) { float res = 0; res += a*a*0.315f + b*0.512f + 0.789f; res += a*a*0.15f + b*0.12f + 0.789f; res += a*a*0.35f + b*0.51f + 0.89f; res += a*a*0.31f + b*0.52f + 0.7f; res += a*a*0.4315f + b*0.512f + 0.4789f; res += a*a*0.515f + b*0.132f + 0.7859f; res += a*a*0.635f + b*0.521f + 0.89f; res += a*a*0.731f + b*0.152f + 0.7f; res += a*a*0.1315f + b*0.512f + 0.789f; res += a*a*0.115f + b*0.12f + 0.789f; res += a*a*0.135f + b*0.51f + 0.89f; res += a*a*0.131f + b*0.52f + 0.7f; res += a*a*0.14315f + b*0.512f + 0.4789f; res += a*a*0.1515f + b*0.132f + 0.7859f; res += a*a*0.1635f + b*0.521f + 0.89f; res += a*a*0.1731f + b*0.152f + 0.7f; return res; }

वास्तव में, इसका बहुत अर्थ नहीं है (और यह स्पष्ट है कि इसे बहुत सरल बनाया जा सकता है), लेकिन शुद्ध गणितीय कार्यों के एक सरल प्रदर्शन के रूप में कार्य करता है। इसमें महत्वपूर्ण बात यह है कि यह एक अलग .cpp फ़ाइल में है और यह बहुत अधिक अंकगणित करता है, लेकिन बाद में अधिक।

भाग दो - प्रारंभिक ओपनसीएल

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

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

तो सबसे पहले हम उपलब्ध प्लेटफार्मों की एक सूची है:

 vector<cl::Platform> platforms; cl::Platform::get(&platforms);

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

बस एक और छोटी सी चाल, C ++ आवरण अपने स्वयं के वैक्टर का उपयोग कर सकता है (यदि आप इसके बारे में उसे बताते हैं) या एसटीडी से वैक्टर, तो अगर कोई ऐसा कुछ है जैसे cl :: वेक्टर उदाहरणों में सामने आता है, तो वह डरो मत, वह दोनों स्वरूपों को जानता है ।

प्लेटफ़ॉर्म की सूची प्राप्त करने के बाद, प्रत्येक प्लेटफ़ॉर्म के लिए हमें उपलब्ध उपकरणों की एक सूची मिलती है:

 std::vector<cl::Device> devices; platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices);

वास्तव में उपकरण वे हैं जो हमारी गणना को अंजाम देंगे। यह एक GPU, एक सीपीयू, या कुछ विशेष त्वरक हो सकता है जो मेजबान से जुड़ा हुआ है, अर्थात्। सिस्टम जिस पर OpenCL चल रहा है। CL_DEVICE_TYPE_ALL के बजाय, आप CL_DEVICE_TYPE_GPU पास कर सकते हैं, फिर यह केवल केंद्रीय प्रोसेसर के लिए वीडियो कार्ड या CL_DEVICE_TYPE_CPU जारी करेगा।

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

भाग तीन - एक कर्नेल बनाना और लॉन्च करना

यहां हम सबसे दिलचस्प भाग पर आते हैं - गणना।

शून्य PerformTestOnDevice (सीएल :: डिवाइस डिवाइस)

 void PerformTestOnDevice(cl::Device device) { cout << endl << "-------------------------------------------------" << endl; cout << "Device: " << device.getInfo<CL_DEVICE_NAME>() << endl << endl; //For the selected device create a context vector<cl::Device> contextDevices; contextDevices.push_back(device); cl::Context context(contextDevices); //For the selected device create a context and command queue cl::CommandQueue queue(context, device); //Clean output buffers fill_n(pOutputVector, DATA_SIZE, 0); //Create memory buffers cl::Buffer clmInputVector1 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector1); cl::Buffer clmInputVector2 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector2); cl::Buffer clmOutputVector = cl::Buffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pOutputVector); //Load OpenCL source code std::ifstream sourceFile("OpenCLFile1.cl"); std::string sourceCode(std::istreambuf_iterator<char>(sourceFile),(std::istreambuf_iterator<char>())); //Build OpenCL program and make the kernel cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1)); cl::Program program = cl::Program(context, source); program.build(contextDevices); cl::Kernel kernel(program, "TestKernel"); //Set arguments to kernel int iArg = 0; kernel.setArg(iArg++, clmInputVector1); kernel.setArg(iArg++, clmInputVector2); kernel.setArg(iArg++, clmOutputVector); kernel.setArg(iArg++, DATA_SIZE); //Some performance measurement timeValues.clear(); __int64 start_count; __int64 end_count; __int64 freq; QueryPerformanceFrequency((LARGE_INTEGER*)&freq); //Run the kernel on specific ND range for(int iTest=0; iTest<TESTS_NUMBER; iTest++) { QueryPerformanceCounter((LARGE_INTEGER*)&start_count); queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(DATA_SIZE), cl::NDRange(128)); queue.finish(); QueryPerformanceCounter((LARGE_INTEGER*)&end_count); double time = 1000 * (double)(end_count - start_count) / (double)freq; timeValues.push_back(time); } PrintTimeStatistic(); // Read buffer C into a local list queue.enqueueReadBuffer(clmOutputVector, CL_TRUE, 0, DATA_SIZE * sizeof(float), pOutputVector); }

सबसे पहले, हम इस तरह से प्राप्त डिवाइस का नाम प्रदर्शित करते हैं:

 device.getInfo<CL_DEVICE_NAME>()

उसी तरह, आप कोर, आवृत्ति, संस्करण, आदि की संख्या के बारे में जानकारी प्राप्त कर सकते हैं।

फिर हम संदर्भ बनाते हैं:

 vector<cl::Device> contextDevices; contextDevices.push_back(device); cl::Context context(contextDevices);

संदर्भों के साथ सब कुछ इतना सरल नहीं है ... एक संदर्भ बनाते समय, हम उन उपकरणों की एक सूची पास करते हैं जिन्हें हम इसमें शामिल करना चाहते हैं, लेकिन एक सीमा है: केवल एक ही मंच पर डिवाइस एक ही संदर्भ में हो सकते हैं, अर्थात GPU और CPU के साथ संदर्भ बनाएं (Intel / NVidia के मामले में) विफल हो जाएगा। एक संदर्भ में कई उपकरणों के मामले में, सभी बफ़र्स को विभिन्न उपकरणों पर स्वचालित रूप से सिंक्रनाइज़ किया जाएगा। एक ओर, यह मल्टी-जीपीयू समर्थन को सरल करता है, और दूसरी ओर, कोई भी नहीं जानता कि ड्राइवर कैसे, क्या और कब सिंक्रनाइज़ करेगा, और उच्च प्रदर्शन प्राप्त करने के लिए डेटा ट्रांसफर दक्षता महत्वपूर्ण है जिसके लिए सब कुछ कल्पना की गई है। इसलिए, मैं आमतौर पर प्रत्येक डिवाइस के लिए एक अलग संदर्भ बनाता हूं और डेटा को मैन्युअल रूप से वितरित करता हूं। इस प्रकार, यह हमेशा ज्ञात होता है कि क्या, कहाँ, कब होता है।

अगला कदम डिवाइस के लिए कमांड कतार बनाना है:

 cl::CommandQueue queue(context, device);

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

अगला, हम इनपुट और आउटपुट वैक्टर के लिए बफ़र्स बनाते हैं:

 //Create memory buffers cl::Buffer clmInputVector1 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector1); cl::Buffer clmInputVector2 = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pInputVector2); cl::Buffer clmOutputVector = cl::Buffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, DATA_SIZE * sizeof(float), pOutputVector);

बफर बनाते समय, संदर्भ (एक विशिष्ट डिवाइस के बजाय) को इंगित किया जाता है, इसकी मात्रा और, यदि वांछित है और CL_MEM_COPY_HOST_PTR ध्वज का उपयोग करते हैं, तो डेटा का एक संकेतक जिसे बनाते समय इसे कॉपी किया जाएगा। जैसा कि मैंने पहले कहा, C ++ रैपर संदर्भ गणना का उपयोग करता है, इसलिए आपको शुद्ध C API के विपरीत, मैन्युअल रूप से बफर को हटाने की आवश्यकता नहीं है।

अगला, हमें एक कर्नेल बनाने की आवश्यकता है जिसका कोड "OpenCLFile1.cl" फ़ाइल में संग्रहीत है। ऐसा करने के लिए, हम फ़ाइल से पाठ पढ़ते हैं, एक ओपनसीएल प्रोग्राम बनाते हैं, इसे संकलित करते हैं, और इसमें से "टेस्टकर्नेल" नाम के साथ कर्नेल प्राप्त करते हैं, जिसे आप अगले भाग में देखेंगे।

 cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1)); cl::Program program = cl::Program(context, source); program.build(contextDevices); cl::Kernel kernel(program, "TestKernel");

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

इसके बाद, हमें उन तर्कों को सेट करने की आवश्यकता है जिन्हें कर्नेल को पास किया जाएगा। CUDA के विपरीत, आपको प्रत्येक तर्क के लिए विशेष फ़ंक्शन (C ++ रैपर, विधियों के मामले में) कॉल करने की आवश्यकता है और यदि आवश्यक हो, तो तर्क का आकार निर्दिष्ट करें।

 int iArg = 0; kernel.setArg(iArg++, clmInputVector1); kernel.setArg(iArg++, clmInputVector2); kernel.setArg(iArg++, clmOutputVector); kernel.setArg(iArg++, DATA_SIZE);

अब हम सबसे महत्वपूर्ण बात पर आते हैं - कर्नेल का प्रक्षेपण:

 queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(DATA_SIZE), cl::NDRange(128));

वास्तव में queue.enqueueNDRangeKernel कमांड की कतार में कर्नेल स्टार्ट कमांड जोड़ता है और संसाधित होने वाले तत्वों की संख्या, साथ ही समूह का आकार निर्धारित करता है। मैं समूहों पर अलग से (दूसरे लेख में) चर्चा करूंगा, लेकिन अब मैं केवल इस तथ्य का उल्लेख करूंगा कि सभी तत्व हमेशा समूहों में विभाजित होते हैं और प्रदर्शन समूह के आकार पर बहुत निर्भर कर सकता है। हमारे मामले में, तत्वों की संख्या DATA_SIZE है, और समूह का आकार 128 है। कर्नेल निष्पादन के दौरान, इसे DATA_SIZE बार (एक अज्ञात अनुक्रम में और संभवतः एक ही समय में) लॉन्च किया जाएगा और प्रत्येक बार लॉन्च होने पर, किस तत्व को संसाधित किया जाता है, इसके बारे में जानकारी भेजी जाएगी।
enqueueNDRangeKernel ब्लॉक नहीं हो रहा है, इसलिए कर्नेल शुरू होने के बाद, हमें इसके समाप्त होने की प्रतीक्षा करनी चाहिए, जिसके लिए यह आवश्यक है:

 queue.finish();

वास्तव में, फिनिश दो कार्य करता है:
1) सभी कमांड को डिवाइस पर भेजता है (enqueueNDRangeK कर्नेल को निष्पादित करना सुनिश्चित करता है कि ड्राइवर कमांड प्राप्त करता है और उसे कतार में खड़ा करता है, लेकिन यह गारंटी नहीं देता है कि यह डिवाइस पर चलेगा, और अक्सर कर्नेल वास्तव में शुरू होने से पहले काफी लंबा समय ले सकता है)।
2) कतार में सभी टीमों के पूरा होने का इंतजार।
यदि आप केवल पहले भाग को निष्पादित करना चाहते हैं, तो एक पुश (clFlush) कमांड है जो अवरुद्ध नहीं है, लेकिन चालक को कतार से कमांड निष्पादित करने के लिए मजबूर करता है।

गणना करने के बाद, हम खर्च किए गए समय की गणना करते हैं और गणना के परिणामों को कमांड के साथ मेजबान पर वापस अपलोड करते हैं:

 queue.enqueueReadBuffer(clmOutputVector, CL_TRUE, 0, DATA_SIZE * sizeof(float), pOutputVector);

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

भाग चार - ओपनसीएल कर्नेल कोड

और यहां हमें वह स्थान मिला, जहां हमें ओपनसीएल पर कोड लिखना शुरू करना होगा (हालाँकि इसे कोड कहना मुश्किल है, इसलिए ... लाड़ करना :))। यह वही है जो OpenCLFile1.cl जैसा दिखता है:

 #include "MathCode.cpp" __kernel void TestKernel( __global const float* pInputVector1, __global const float* pInputVector2, __global float* pOutputVectorHost, int elementsNumber) { //Get index into global data array int iJob = get_global_id(0); //Check boundary conditions if (iJob >= elementsNumber) return; //Perform calculations pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]); }

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

 int iJob = get_global_id(0);

Get_global_id पैरामीटर एक आयाम को इंगित करता है, क्योंकि संसाधित किए जा रहे तत्व 1, 2 या 3-आयामी सरणी हो सकते हैं।
फिर हम सीमा की स्थिति की जाँच करते हैं:

 if (iJob >= elementsNumber) return;

यह इस कारण से किया जाना चाहिए कि संसाधित किए जाने वाले तत्वों की संख्या हमेशा समूह के आकार के एक से अधिक होनी चाहिए और इस प्रकार यह संसाधित होने वाली संख्या से अधिक हो सकती है।
और जाँच के बाद, हम मुख्य भाग करते हैं: गणना, ठीक उसी तरह जैसे कि मेजबान पर:

 pOutputVectorHost[iJob] = MathCalculations(pInputVector1[iJob], pInputVector2[iJob]);

भाग पांच - परीक्षण और माप प्रदर्शन

तो यह एप्लिकेशन लॉन्च करने, प्रदर्शन का मूल्यांकन करने और कुछ निष्कर्ष निकालने का समय है।

मैंने दो मशीनों पर परीक्षण चलाया और दिलचस्प परिणाम मिले:
नोटबुक (CPU: Intel® Core ™ i7-820QM , GPU: NVidia Quadro FX 2800M ):

 Host: 959.256 ms CPU: 82.4163 ms (13.106X faster than host) GPU: 9.90836 ms (109.014X faster than host)

डेस्कटॉप (CPU: Intel® Core ™ i7-2600 , GPU: NVidia GeForce GTX 580 )

 Host: 699.031 ms CPU: 27.7833 ms (25.159X faster than host) GPU: 2.06257 ms (338.897X faster than host)

पूर्ण परिणाम

 Device: Host Calculation time statistic: (20 runs) Med: 959.256 ms (1.12602X faster than host) Avg: 1080.15 ms Min: 933.554 ms Max: 1319.19 ms ------------------------------------------------- Device: Quadro FX 2800M Calculation time statistic: (200 runs) Med: 9.90836 ms (109.014X faster than host) Avg: 10.7231 ms Min: 9.82841 ms Max: 135.924 ms Errors: avgRelAbsDiff = 5.25777e-008 maxRelAbsDiff = 5.83678e-007 ------------------------------------------------- Device: Intel(R) Core(TM) i7 CPU Q 820 @ 1.73GHz Calculation time statistic: (200 runs) Med: 82.4163 ms (13.106X faster than host) Avg: 85.2226 ms Min: 79.4138 ms Max: 113.03 ms Errors: avgRelAbsDiff = 3.64332e-008 maxRelAbsDiff = 4.84797e-007

 Device: Host Calculation time statistic: (20 runs) Med: 699.031 ms (0.999956X faster than host) Avg: 699.1 ms Min: 691.544 ms Max: 715.233 ms ------------------------------------------------- Device: GeForce GTX 580 Calculation time statistic: (200 runs) Med: 2.06257 ms (338.897X faster than host) Avg: 2.4 ms Min: 2.03873 ms Max: 82.0514 ms Errors: avgRelAbsDiff = 3.50006e-008 maxRelAbsDiff = 4.92271e-007 ------------------------------------------------- Device: Intel(R) Core(TM) i7-2600 CPU @ 3.40GHz Calculation time statistic: (200 runs) Med: 27.7833 ms (25.159X faster than host) Avg: 27.49 ms Min: 27.0154 ms Max: 35.8386 ms Errors: avgRelAbsDiff = 3.64377e-008 maxRelAbsDiff = 4.89584e-007

तो, चलो परिणामों का विश्लेषण शुरू करते हैं, और परिणाम, मुझे कहना होगा, बहुत प्रभावशाली हैं। लैपटॉप पर GPU होस्ट की तुलना में ~ 110X तेज है, और डेस्कटॉप पर यह ~ 340X तेज, एक प्रभावशाली परिणाम है। इससे पहले कि आप मुझ पर चप्पल फेंकना शुरू करें और कहें कि इस तरह की तुलना सही नहीं है, मैं कहूंगा कि वास्तव में इसमें कुछ ट्रिक्स हैं, लेकिन इससे ज्यादा कुछ नहीं।

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

दूसरे, याद रखें कि गणित कोड कैसा दिखता था? जिन लोगों ने इसे नहीं देखा, उनके लिए मैं कहूंगा कि यह एक ही डेटा पर बहुत सारे गणितीय ऑपरेशन हैं, और यह सरल कॉपी-पेस्ट से निकला और गुणांक में संख्याओं की जगह ले रहा है, लेकिन शुरू में यह सरल था और केवल एक पंक्ति में ले गया, केवल जब मैंने परीक्षण करना शुरू किया, परिणाम इतने खुश नहीं थे, GPU केवल 4-5 गुना तेज था। आप ऐसा क्यों सोचते हैं? (एक लफ्फाजी वाला सवाल, आप सोच भी नहीं सकते :))। और सब कुछ सरल है, हम स्मृति प्रदर्शन में भाग गए। मुझे उम्मीद है कि बाद में मुझे इस पर हाथ मिलेगा और मैं मेमोरी और प्रोसेसर के प्रदर्शन के बीच संबंध पर एक लेख लिखूंगा, लेकिन यह एक अलग कहानी है, इस लेख में हम केवल इस तथ्य में रुचि रखते हैं कि इस कर्नेल के साथ हमें अंकगणितीय प्रोसेसर प्रदर्शन का एक स्वच्छ परीक्षण मिला।

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

लेकिन हम जानते हैं कि GPU जल्दी से इस पर विचार करता है, और हमारे परीक्षण के परिणामस्वरूप, यह पता चला कि सीपीयू OpenCL कोड को जल्दी से निष्पादित करता है, सटीक होने के लिए, यह डिफ़ॉल्ट सेटिंग्स के साथ MSVC10 द्वारा संकलित नियमित कोड की तुलना में 13X और 25X तेज है। आइए समझते हैं कि यह कैसे निकलता है और ये नंबर कहां से आते हैं।

दोनों प्रोसेसर में 4 वास्तविक और 8 वर्चुअल कोर हैं, और ओपनसीएल सिर्फ सभी कोर का उपयोग करने के लिए बनाया गया है, लेकिन हमारे पास 4X के लिए बहुत अधिक सुधार है। और यहां मुझे इंटेल के लिए धन्यवाद कहना चाहिए, जिसने ओपनसीएल के कार्यान्वयन में, स्वचालित वेक्टराइजेशन के लिए समर्थन जोड़ा, अर्थात। बिना किसी कोड परिवर्तन के, OpenCL SSE या AVX का उपयोग करता है, जो कि उपलब्ध है पर निर्भर करता है। यह देखते हुए कि हमारे पास 128-बिट SSE है, और AVX 256-बिट के साथ काम करता है, यह पता चलता है कि प्रदर्शन क्रमशः 16X और 32X तक बढ़ना चाहिए। यह सच्चाई के करीब है, लेकिन अभी भी बिल्कुल सटीक संयोग नहीं है। और फिर हमें TurboBoost जैसी आनंददायक चीज़ को याद रखने की आवश्यकता है। ये प्रोसेसर 1.73 गीगाहर्ट्ज / 3.06 गीगाहर्ट्ज (लैपटॉप) और 3.4 गीगाहर्ट्ज / 3.8 गीगाहर्ट्ज (डेस्कटॉप) की आवृत्ति पर काम करते हैं, लेकिन वास्तव में मैं कह सकता हूं कि लैपटॉप प्रोसेसर की आवृत्ति 1.73 से 2.8 तक लगातार कूदती है, और यह बहुत दृढ़ता से गर्म होता है (यहां आपको घुमावदार शीतलन प्रणाली के लिए डेल पर एक बड़ा कैमियो फेंकना चाहिए), इसलिए 3.06GHz आवृत्ति परीक्षण के दौरान हमने वास्तव में कोई महत्वपूर्ण समय नहीं देखा है। साथ ही, आपको यह नहीं भूलना चाहिए कि व्यावहारिक परिणाम हमेशा सैद्धांतिक रूप से संभव से कम है (डेस्कटॉप को तेजी से काम करना चाहिए, सिद्धांत रूप में), लेकिन जैसा कि हम देख सकते हैं, 25X प्रदर्शन में सुधार एक ही हार्डवेयर पर लगभग मुफ्त में प्राप्त किया जा सकता है।

निष्कर्ष

इस लेख का लक्ष्य OpenCL के साथ काम करने के सभी विवरणों को समझाने का प्रयास नहीं था, बल्कि यह दिखाने का प्रयास था कि सब कुछ इतना जटिल नहीं है ( यहाँ मैंने पहले ही लिखा था कि यह इतना सरल नहीं है) और आदर्श परिस्थितियों में आप बहुत प्रभावशाली प्रदर्शन प्राप्त कर सकते हैं, और यहां तक कि एक ही हार्डवेयर पर, और इसके अलावा, आप सभी उपकरणों के लिए एक ही कोड का उपयोग कर सकते हैं। लेकिन याद रखें कि ये लगभग आदर्श स्थितियां हैं, जो हमेशा नहीं होती हैं।

पुनश्च: उन लोगों के लिए जो कोड के साथ खेलना चाहते हैं और किसी अन्य हार्डवेयर पर परीक्षण देखते हैं, परियोजना (और यहां तक कि इकट्ठे निष्पादन योग्य) गीथब पर स्थित है । शुरू करने के लिए, आपको अपने हार्डवेयर के निर्माताओं से OpenCL SDK की आवश्यकता हो सकती है।

PS2: यदि किसी के पास एक आइवी ब्रिज है, तो एकीकृत वीडियो कोर के परीक्षण को देखना दिलचस्प होगा। तथ्य यह है कि ओपनसीएल एसडीके के नवीनतम संस्करण में, इंटेल ने आईजीपी तक पहुंच खोली है, लेकिन केवल नवीनतम पीढ़ी के प्रोसेसर के लिए, लेकिन मेरे पास ऐसा नहीं है। और एएमडी के परिणाम देखने में दिलचस्प हैं।

OpenCL के साथ पहला कदम या GPU और CPU पर समान कोड को चलाने की कहानी