CUDA का उपयोग करके स्टैंसिल बफर कार्यान्वयन उदाहरण

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


स्टैंसिल को गहराई के साथ एक बफर में संग्रहीत किया जाता है। उदाहरण के लिए, डी 3 डीएफएमटी_ डी 24 एस 8 प्रारूप में, 24 बिट गहराई बिट और 8 बिट स्टेंसिल हैं। सादगी के लिए, हम आगे मानेंगे कि स्टैंसिल अंतिम बिट में संग्रहीत है। यदि यह बिट = 1 है, तो पिक्सेल सक्रिय है। इस प्रकार, सरलीकृत एल्गोरिथ्म निम्नलिखित कदम प्रस्तुत करता है:

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

इस प्रकार, जहां दर्पण छवि में था, प्रतिबिंब दिखाई देगा। और जहां कोई नहीं है, वहां कुछ भी नहीं बदलेगा।

CUDA सॉफ्टवेयर कार्यान्वयन

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

तो, हम स्टैंसिल बफर को (N / 32) * sizeof (int) बाइट्स के आकार से शुरू करते हैं। और इसकी बनावट को बांध दें।

cudaMalloc((void**)&m_stencilBuffer, N*sizeof(int)/32); cudaBindTexture(0, stencil_tex, m_stencilBuffer, N*sizeof(int)/32);     -  (.h )  : Texture<int, 1, cudaReadModeElementType> stencil_tex; ,        : static __device__ int g_stencilMask[32] = { 0x00000001, 0x00000002, 0x00000004, 0x00000008, 0x00000010, 0x00000020, 0x00000040, 0x00000080, 0x00000100, 0x00000200, 0x00000400, 0x00000800, 0x00001000, 0x00002000, 0x00004000, 0x00008000, 0x00010000, 0x00020000, 0x00040000, 0x00080000, 0x00100000, 0x00200000, 0x00400000, 0x00800000, 0x01000000, 0x02000000, 0x04000000, 0x08000000, 0x10000000, 0x20000000, 0x40000000, 0x80000000 }; 

उन गुठली के लिए जो केवल स्टैंसिल बफर को पढ़ते हैं, कर्नेल की शुरुआत में मैक्रो का उपयोग निम्नानुसार करते हैं:

 __global__ void my_kernel(…) { uint tid = blockDim.x * blockIdx.x + threadIdx.x; STENCIL_TEST(tid); // my code here } 

व्यवहार में (GTX560), इस तरह के स्टैंसिल परीक्षण एक साधारण जांच की तुलना में लगभग 20-25% तेज है, प्रकार की जांच:

 uint activeFlag = a_flags[tid]; if(activeFlag==0) return; 

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

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

 // (tid >> 5) same as (tid/32) // (tid & 0x1f) same as (tid%32) __global__ void my_kernel2(…,uint* a_stencilBuffer) { uint tid = blockDim.x * blockIdx.x + threadIdx.x; uint activeWarp = a_stencilBuffer[tid >> 5]; if(activeWarp==0) // all threads in warp inactive return; // each threads will store it's particular bit from group of 32 threads uint active = activeWarp & g_stencilMask[tid&0x1f]; if(!active) goto WRITE_BACK_STENCIL_DATA; // my code here WRITE_BACK_STENCIL_DATA: WriteStencilBit(tid, a_stencilBuffer, active); } 

यदि धागा निष्क्रिय है, तो यह तुरंत कर्नेल स्केट में जाएगा। यदि किसी कारण से आप अपने कोड के अंदर यह तय करते हैं कि यह धागा निष्क्रिय होना चाहिए, तो यह करें:

 if(want to kill thread) { active = 0; goto WRITE_BACK_STENCIL_DATA; } 

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

 __device__ void WriteStencilBit(int tid, uint* a_stencilBuffer, uint value) { uint stencilMask = __ballot(value); if((tid & 0x1f) == 0) // same as tid%32 == 0 a_stencilBuffer[tid >> 5] = stencilMask; } 

__Ballot () फ़ंक्शन uint देता है, जहां प्रत्येक ith बिट 1 है यदि और केवल यदि इसके तर्क में क्या है तो शून्य नहीं है। यही है, यह वही करता है जो इसे वहां की जरूरत है, झंडे को वापस ताना के अंदर अलग-अलग धागे से उंट में सिलाई करना।

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

 __device__ void WriteWrongStencilBit(int tid, uint* a_stencilBuffer, uint value) { if((tid & 0x1f) == 0) // same as tid%32 == 0 a_stencilBuffer[tid >> 5] = __ballot(value); } 

तथ्य यह है कि __ballot () हमेशा उन बिट्स पर 0 लगाएगा जिनकी धाराएं वर्तमान में नकाबपोश हैं। और नॉन-ज़ीरो (1..31) के अंदर एक नंबर के साथ सभी स्ट्रीम नकाबपोश होंगे और अगर स्टेटमेंट के अंदर नहीं पहुंचेंगे, इसलिए इस तरह के कोड के लिए __ballot () फ़ंक्शन के 1..31 बिट्स हमेशा शून्य होंगे। यहाँ से सत्य एक दिलचस्प निष्कर्ष का अनुसरण करता है। यदि आपको फर्मी वास्तुकला के साथ वीडियो कार्ड के लिए लिखने की गारंटी दी जाती है, तो यहां तक ​​कि गुठली के लिए जो स्टैंसिल बफर को लिखते हैं, आप निम्नानुसार धारा को मार सकते हैं:

 if(want to kill thread) return; 

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

 if(want to kill thread && (tid&0x1f!=0)) return; 

या ऊपर दिए गए फॉर्म का उपयोग करें:

 if(want to kill thread) { active = 0; goto WRITE_BACK_STENCIL_DATA; } 

पुराने हार्डवेयर के लिए कार्यान्वयन सुविधाएँ (G80-GT200)

आइए अब विचार करें कि स्टैंसिल के लिए पुराने जीपीयू पर प्रभावी ढंग से काम करने के लिए क्या एक्सटेंशन होना चाहिए। इन वीडियो कार्ड पर __ballot () फ़ंक्शन समर्थित नहीं है। हम अपने पास मौजूद क्षमताओं के अनुसार WriteStencilBit फ़ंक्शन को फिर से लिखते हैं:

 template<int CURR_BLOCK_SIZE> __device__ inline void WriteStencilBit(int tid, uint* a_stencilBuffer, uint value) { #if COMPUTE_CAPABILITY >= COMPUTE_CAPABILITY_GF100 uint stencilMask = __ballot(value); if((tid & 0x1f) == 0) a_stencilBuffer[tid >> 5] = stencilMask; #elif COMPUTE_CAPABILITY >= COMPUTE_CAPABILITY_GT200 if(__all(value==0)) { if((tid & 0x1f) == 0) a_stencilBuffer[tid >> 5] = 0; } else if(__all(value)) { if((tid & 0x1f) == 0) a_stencilBuffer[tid >> 5] = 0xffffffff; } else { __shared__ uint active_threads[CURR_BLOCK_SIZE/32]; uint* pAddr = active_threads + (threadIdx.x >> 5); if((tid & 0x1f) == 0) *pAddr = 0; atomicOr(pAddr, value); if((tid & 0x1f) == 0) a_stencilBuffer[tid >> 5] = *pAddr; } #else __shared__ uint active_threads[CURR_BLOCK_SIZE]; active_threads[threadIdx.x] = value; if((threadIdx.x & 0x1) == 0) active_threads[threadIdx.x] = value | active_threads[threadIdx.x+1]; if((threadIdx.x & 0x3) == 0) active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+2]; if((threadIdx.x & 0x7) == 0) active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+4]; if((threadIdx.x & 0xf) == 0) active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+8]; if((threadIdx.x & 0x1f) == 0) active_threads[threadIdx.x] = active_threads[threadIdx.x] | active_threads[threadIdx.x+16]; uint* perWarpArray = active_threads + ((threadIdx.x >> 5) << 5); if((tid & 0x1f) == 0) a_stencilBuffer[tid >> 5] = perWarpArray[0]; #endif } 

इस प्रकार हम साझा मेमोरी + 2 मतदान कार्यों में परमाणु बना सकते हैं, __any और __all उपलब्ध हैं, ताकि हम उनका उपयोग कर सकें। अन्य मामलों में, केवल शास्त्रीय कमी बनी हुई है।

स्टैंसिल का परीक्षण

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

एफपीएस की गतिशीलता इस प्रकार है: 30, 25, 23.7, 20, 19.4, 18.8