CUDA का उपयोग करके एक विद्युत क्षेत्र की मॉडलिंग करना

यह लेख यह प्रदर्शित करने के लिए लिखा गया है कि कैसे CUDA तकनीक आवेशित कणों की सरल अंतःक्रिया को अनुकरण कर सकती है (देखें Coulomb’s Law ) एक स्थिर चित्र प्रदर्शित करने के लिए, मैंने फ्रीग्लूट लाइब्रेरी का उपयोग किया।
जैसा कि वे अक्सर हैबे पर लिखते हैं:

टर्बो पास्कल और बोरलैंड सी 5.02 के दिनों में, मुझे प्रकृति की सरल भौतिक घटनाओं के अनुकरण की प्रक्रिया में दिलचस्पी थी - चाहे वह गुरुत्वाकर्षण आकर्षण बल हो या कूलम्ब इंटरैक्शन। दुर्भाग्य से, कई सौ मेगाहर्ट्ज वाले एक प्रोसेसर की शक्ति सरल गणितीय कार्यों के लिए पर्याप्त नहीं थी, जिनमें से कई थे।

यद्यपि मैं एक पोग्रोमिस्ट नहीं हूं, लेकिन मैं एक बार एक दूर के बचपन को याद करना चाहता था, और जावास्क्रिप्ट की मूल बातें सीखना चाहता था, और किसी तरह मैंने किसी भी उचित संख्या में सामग्री बिंदुओं के गुरुत्वाकर्षण बातचीत का अनुकरण किया। इसलिए, उदाहरण के लिए, न्यूटन के नियम का उपयोग करके नीचे का कण "सूर्य" के चारों ओर घूमता है:

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

तो, उदासीन बंद करो, चलो pogrom कार्यक्रम। मेरा मंच उबंटू वाला एक लैपटॉप है और एक GeForce GT 540m ग्राफिक्स कार्ड और CUDA टूलकिट 5.0 स्थापित है। यदि आप CUDA के साथ शुरुआत कर रहे हैं, तो मैं आपको एक दिलचस्प पुस्तक, CUDA by Example पर सलाह दे सकता हूं।

सबसे पहले, Nsight टेम्प्लेट से एक प्रोजेक्ट बनाएं और इसे GLU, GL, ग्लूट लाइब्रेरीज़ का उपयोग करके फ्रीग्लूट से कनेक्ट करें। यह हमें परिणामी वेक्टर फ़ील्ड प्रदर्शित करने की अनुमति देगा, जिसकी हम गणना करते हैं। main फ़ंक्शन के साथ हमारी फ़ाइल की अनुमानित संरचना निम्नानुसार होगी:

 #include <stdio.h> #include <stdlib.h> int main(int argc, char** argv) { // Initialize freeglut // ... // Initialize Host and Device variables // ... // Launch Kernel to render the image // Display Image // Free resources return 0; } 

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

विंडो आरंभीकरण

जब से मैंने आवेदन की कल्पना की, मैंने खिड़की को गतिशील रूप से आकार देने की योजना नहीं बनाई थी, मैंने वैश्विक वैज्ञानिको का उपयोग करते हुए इसे घोषित करने का निर्णय लिया:

 const int width = 1280; const int height = 720; 

अरे हाँ! C ++ और C99 शैलियों के मिश्रण के लिए कड़ाई से न्याय न करें। मुझे कभी-कभी कुछ चीजें मिल जाती हैं, जैसे कि ओवरलोडिंग, उपयोगी, लेकिन मैंने किसी तरह से सीयूडीए के साथ अनुप्रयोगों की शैली का पालन करने के लिए मॉलॉक को नहीं छोड़ा है।

विंडो ही, एस्केप कुंजी को प्रोसेस करना और सरल कोड निम्नलिखित कोड के साथ किया जा सकता है:

 #include <GL/freeglut.h> #include <GL/gl.h> #include <GL/glext.h> // ... void key(unsigned char key, int x, int y) { switch (key) { case 27: printf("Exit application\n"); glutLeaveMainLoop(); break; } } void draw(void) { glClearColor(0.0, 0.0, 0.0, 1.0); glClear(GL_COLOR_BUFFER_BIT); // glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, screen); glFlush(); } // ... // Initialize freeglut glutInit(&argc, argv); glutInitDisplayMode(GLUT_SINGLE | GLUT_RGBA); glutInitWindowSize(width, height); glutCreateWindow("Electric field"); glutDisplayFunc(draw); glutKeyboardFunc(key); glutSetOption(GLUT_ACTION_ON_WINDOW_CLOSE, GLUT_ACTION_CONTINUE_EXECUTION); // ... // Display Image glutMainLoop(); 

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

• glutLeaveMainLoop - जब एस्केप दबाया जाता है, तो फ्रीग्लुट आपको मुख्य लूप को बाधित करने की अनुमति देता है ताकि हम एप्लिकेशन संसाधनों को मुक्त कर सकें। यदि आप चाहते हैं, तो आप atexit उपयोग कर सकते हैं;
• glDrawPixels - बाद में मैं unsigned byte से एक बफर glDrawPixels , जहां मैं सभी पिक्सेल (उनके रंग) संग्रहीत करूंगा और फिर उन्हें एक विंडो में प्रदर्शित करूंगा;
• glutSetOption - यह मुझे गतिशील मेमोरी के रूप में main से स्पष्ट संसाधनों में कोड निष्पादित करना जारी रखने की अनुमति देता है।

ब्लॉक और थ्रेड में विभाजन की गणना

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

धाराएं, बदले में, एक आयामी पट्टी या दो-आयामी तालिका के रूप में एक ब्लॉक में भी स्थित हो सकती हैं। वैसे, तीन आयामों के साथ विकल्प संभव है, लेकिन मैं उन्हें याद करूंगा, क्योंकि वे अभी तक मेरे कार्य में फिट नहीं हैं।

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

किसी ब्लॉक के अंदर थ्रेड्स की संख्या जानने के लिए, हम cudaGetDeviceProperties फ़ंक्शन का उपयोग करते हैं:

 int threadsCount(void) { int deviceCount = 0; cudaGetDeviceCount(&deviceCount); if (deviceCount <= 0) { printf("No CUDA devices\n"); exit(-1); } cudaDeviceProp properties; cudaGetDeviceProperties(&properties, 0); return properties.maxThreadsPerBlock; } 

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

अब हम अपने क्षेत्र को ब्लॉक और प्रवाह में तोड़ देंगे:

 void setupGrid(dim3* blocks, dim3* threads, int maxThreads) { threads->x = 32; threads->y = maxThreads / threads->x - 2; // to avoid cudaErrorLaunchOutOfResources error blocks->x = (width + threads->x - 1) / threads->x; blocks->y = (height + threads->y - 1) / threads->y; } // ... // Initialize Host dim3 blocks, threads; setupGrid(&blocks, &threads, threadsCount()); printf("Debug: blocks(%d, %d), threads(%d, %d)\nCalculated Resolution: %dx %d\n", blocks.x, blocks.y, threads.x, threads.y, blocks.x * threads.x, blocks.y * threads.y); 

परिणाम निम्नानुसार होगा (मैंने जानबूझकर उन ब्लॉकों की संख्या में एक मामले में वृद्धि की है जहां स्क्रीन का आकार थ्रेड्स की संख्या का एक से अधिक नहीं है):

 Debug: blocks(40, 24), threads(32, 30) Calculated Resolution: 1280 x 720 

सिद्धांत और अभ्यास

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

एक शुल्क निर्देशांक के साथ एक संरचना है (जैसे स्क्रीन पर)।

 struct charge { int x, y, q; }; 

एक बिंदु पर विद्युत क्षेत्र की ताकत एक वेक्टर है जिसे हम अनुमानों और फ़ंक्शन द्वारा व्यक्त करते हैं जो उन्हें गणना करता है (यह वह जगह है जहां मुझे कोई त्रुटि मिली:

 const float k = 50.0f; const float minDistance = 0.9f; // not to divide by zero const float maxForce = 1e7f; struct force { float fx, fy; __device__ force() : fx(0.0f), fy(0.0f) { } __device__ force(int fx, int fy) : fx(fx), fy(fy) { } __device__ float length2() const { return (fx * fx + fy * fy); } __device__ float length() const { return sqrtf(length2()); } __device__ void calculate(const charge& q, int probe_x, int probe_y) { // F(1->2) = k * q1 * q2 / r(1->2)^2 * vec_r(1->2) / abs(vec_r(1->2)) // e = vec_F / q2 fx = probe_x - qx; fy = probe_y - qy; float l = length(); if (l <= minDistance) { return; } float e = k * qq / (l * l * l); if (e > maxForce) { fx = fy = maxForce; } else { fx *= e; fy *= e; } } __device__ force operator +(const force& f) const { return force(fx + f.fx, fy + f.fy); } __device__ force operator -(const force& f) const { return force(fx - f.fx, fy - f.fy); } __device__ force& operator +=(const force& f) { fx += f.fx; fy += f.fy; return *this; } __device__ force& operator -=(const force& f) { fx -= f.fx; fy -= f.fy; return *this; } }; 

मेरा विचार सरल है - मैं तनाव वाले वैक्टरों का एक दो-आयामी सरणी बनाऊंगा। स्क्रीन पर प्रत्येक बिंदु पर, तनाव कूलम्ब बलों के वैक्टर के योग के बराबर होगा। यह बहुत बुरा लगता है, लेकिन मैं अन्यथा नहीं बना सकता। यह कहना आसान है, कोड इस तरह होगा (आप इसे पढ़ते हैं, लेकिन अपनी परियोजना में अभी तक नहीं लिखते हैं):

  force temp_f; for (int i = 0; i < chargeCount; i++) { temp_f.calculate(charges[i], x, y); *f += temp_f; } 

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

अब जब हमने इन्फ्रास्ट्रक्चर बनाया है - CUDA के लिए कोड। सबसे पहले, आवश्यक होस्ट और डिवाइस चर बनाएं और उन्हें मुक्त करें!

हम संसाधन बनाते हैं और मुक्त करते हैं

हमें निम्नलिखित चर की आवश्यकता होगी:

• स्क्रीन - मेजबान - रंग पिक्सेल;
• स्क्रीन - डिवाइस - वे हैं;
• शुल्क - डिवाइस - डिवाइस की स्थायी मेमोरी में निर्देशांक और परिमाणों का परिमाण।

हम बनाते हैं:

 uchar4* screen = NULL; // ... screen = (uchar4*) malloc(width * height * sizeof(uchar4)); memset(screen, 0, width * height * sizeof(uchar4)); uchar4 *dev_screen = NULL; cudaMalloc((void**) &dev_screen, width * height * sizeof(uchar4)); cudaMemset(dev_screen, 0, width * height * sizeof(uchar4)); // ... // Free resources free(screen); screen = NULL; cudaFree(dev_screen); dev_screen = NULL; 

हम केवल आरोपों को शुरू करते हैं। मैं समान चित्र प्राप्त करने के लिए यादृच्छिक संख्या जनरेटर को रीसेट नहीं करता हूं। यदि आप हर बार एक यादृच्छिक क्षेत्र चाहते हैं, तो आप एक srand कॉल जोड़ सकते हैं:

 const int chargeCount = 10; __constant__ charge dev_charges[chargeCount]; const int maxCharge = 1000; const int minCharge = -1000; // ... void prepareCharges(void) { charge* charges = (charge*) malloc(chargeCount * sizeof(charge)); for (int i = 0; i < chargeCount; i++) { charges[i].x = rand() % width; charges[i].y = rand() % height; charges[i].q = rand() % (maxCharge - minCharge) + minCharge; printf("Debug: Charge #%d (%d, %d, %d)\n", i, charges[i].x, charges[i].y, charges[i].q); } cudaMemcpyToSymbol(dev_charges, charges, chargeCount * sizeof(charge)); } // ... prepareCharges(); 

याद रखें कमेंट आउट लाइन स्क्रीन प्रदान करता है? अब आप इसे अनसुना कर सकते हैं:

  glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, screen); 

एक कर्नेल बनाएँ

अब एक फ़ंक्शन बनाएं जिसे हम प्रत्येक थ्रेड में स्क्रीन पर प्रत्येक पिक्सेल को प्रस्तुत करने के लिए कहेंगे। यह फ़ंक्शन निम्न कार्य करेगा: यह क्षेत्र की ताकत के मूल्य की गणना करेगा, और एक बिंदु पर कुछ रंग से मेल खाने के लिए दूसरे फ़ंक्शन को कॉल करेगा।

 __device__ uchar4 getColor(const force& f) { uchar4 color; color.x = color.y = color.z = color.w = 0; float l = f.length(); color.x = (l > maxLengthForColor ? 255 : l * 256 / maxLengthForColor); return color; } __global__ void renderFrame(uchar4* screen) { // build the field and render the frame } 

कोर को कॉल करने के लिए ... और अभी तक सब कुछ स्पष्ट है:

  // Launch Kernel to render the image renderFrame<<<blocks, threads>>>(dev_screen); cudaMemcpy(screen, dev_screen, width * height * sizeof(uchar4), cudaMemcpyDeviceToHost); 

अब renderFrame कुछ कोड जोड़ें और हमें जो चाहिए वह प्राप्त करें:

 __global__ void renderFrame(uchar4* screen) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; force f, temp_f; for (int i = 0; i < chargeCount; i++) { temp_f.calculate(dev_charges[i], x, y); f += temp_f; } screen[x + y * width] = getColor(f); } 

निष्कर्ष

अक्सर लोग एक लेख को छोड़ देते हैं क्योंकि निष्कर्ष महत्वपूर्ण है। मेरे पास यहां कहने के लिए कुछ भी नहीं है, सिवाय इसके कि मैंने क्या किया। आप यहां लगभग खाली भंडार का लिंक पा सकते हैं।

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