2 संपादित करें: कृपया एक नज़र डालें यह क्रॉसपोस्ट है TLDR के लिए।
संपादित करें: यह देखते हुए कि कणों को ग्रिड कोशिकाओं में विभाजित किया जाता है (कहते हैं 16^3 ग्रिड), क्या प्रत्येक ग्रिड सेल के लिए एक कार्य-समूह और एक कार्य-समूह में कई कार्य-आइटमों को चलाने देना बेहतर है क्योंकि प्रति ग्रिड सेल में कणों की अधिकतम संख्या हो सकती है?
उस स्थिति में मैं सभी कणों को लोड कर सकता हूंस्थानीय कोशिकाओं में पड़ोसी कोशिकाओं और उनके माध्यम से पुनरावृति कुछ गुणों की गणना। तब मैं वर्तमान ग्रिड सेल में प्रत्येक कण में विशिष्ट मूल्य लिख सकता था।
क्या यह दृष्टिकोण सभी कणों के लिए कर्नेल को चलाने और प्रत्येक पर चलने वाले (अधिकांश समय समान) पड़ोसियों के लिए फायदेमंद होगा?
साथ ही, इसका आदर्श अनुपात क्या है number of particles/number of grid cells?
मैं पुन: लागू करने की कोशिश कर रहा हूँ (और संशोधित) CUDA कण OpenCL के लिए और हर कण के लिए निकटतम पड़ोसियों को क्वेरी करने के लिए इसका उपयोग करें। मैंने निम्नलिखित संरचनाएँ बनाई हैं:
- बफर
Pसभी कण पकड़े हुए "3 डी स्थिति (float3) बफर
Spभंडारणint2कण आईडी और उनके स्थानिक हैश की जोड़ी।Spहैश के अनुसार क्रमबद्ध किया जाता है। (हैश सिर्फ 3 डी से 1 डी के लिए एक सरल रैखिक मानचित्रण है - अभी तक कोई जेड-इंडेक्सिंग नहीं है।)बफर
Lभंडारणint2बफर में विशेष स्थानिक हैश की शुरुआत और समाप्ति की स्थिति के जोड़ेSp। उदाहरण:L[12] = (int2)(0, 50).L[12].xसूचकांक है (मेंSp) का प्रथम स्थानिक हैश के साथ कण12.L[12].yसूचकांक है (मेंSp) का अंतिम स्थानिक हैश के साथ कण12.
अब जब मेरे पास ये सभी बफ़र्स हैं, तो मैं सभी कणों के माध्यम से पुनरावृति करना चाहता हूं P और अपने निकटतम पड़ोसियों के माध्यम से प्रत्येक कण पुनरावृति के लिए। वर्तमान में मेरे पास एक कर्नेल है जो इस तरह दिखता है (स्यूडोकोड):
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
for(int x=-1; x<=1; x++)
for(int y=-1; y<=1; y++)
for(int z=-1; z<=1; z++) {
float3 neigh_position = curr_particle + (float3)(x,y,z)*GRID_CELL_SIDE;
// ugly boundary checking
if ( dot(neigh_position<0, (float3)(1)) +
dot(neigh_position>BOUNDARY, (float3)(1)) != 0)
continue;
int neigh_hash = spatial_hash( neigh_position );
int2 particles_range = L[ neigh_hash ];
for(int p=particles_range.x; p<particles_range.y; p++)
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
उस कोड के साथ समस्या यह है कि यह धीमी है। मुझे नॉनलाइनियर जीपीयू मेमोरी एक्सेस (पार्टिसिलरी) पर संदेह है P[Sp[p].y] भीतर-भीतर for पाश) सुस्ती का कारण होना।
मैं जो करना चाहता हूं उसका उपयोग करना है जेड-ऑर्डर वक्र स्थानिक हैश के रूप में। इस तरह मैं केवल 1 कर सकता था for पड़ोसी को क्वेरी करते समय स्मृति की निरंतर श्रेणी के माध्यम से चलने वाला लूप। एकमात्र समस्या यह है कि मुझे पता नहीं है कि क्या होना चाहिए और जेड-इंडेक्स मूल्यों को रोकना चाहिए।
पवित्र कब्र जिसे मैं प्राप्त करना चाहता हूं:
__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid = get_global_id(0);
float3 curr_particle = P[gid];
int processed_value = 0;
// How to accomplish this??
// `get_neighbors_range()` returns start and end Z-index values
// representing the start and end near neighbors cells range
int2 nearest_neighboring_cells_range = get_neighbors_range(curr_particle);
int first_particle_id = L[ nearest_neighboring_cells_range.x ].x;
int last_particle_id = L[ nearest_neighboring_cells_range.y ].y;
for(int p=first_particle_id; p<=last_particle_id; p++) {
processed_value += heavy_computation( P[ Sp[p].y ] );
}
Out[gid] = processed_value;
}
उत्तर:
उत्तर № 1 के लिए -1आपको मॉर्टन कोड एल्गोरिदम का बारीकी से अध्ययन करना चाहिए। एरिक्सन रियल टाइम टक्कर का पता लगाने के बारे में बताता है कि बहुत अच्छी तरह से।
उपयोगकर्ता - वास्तविक समय टकराव का पता लगाने
यहाँ कुछ परीक्षणों सहित एक और अच्छी व्याख्या है:
मॉर्टन एन्कोडिंग / बिट कोडिंग के माध्यम से डिकोडिंग: कार्यान्वयन
जेड-ऑर्डर एल्गोरिदम केवल के रास्तों को परिभाषित करता हैनिर्देशांक जिसमें आप 2 या 3 डी से हैश कर सकते हैं सिर्फ एक पूर्णांक के लिए। यद्यपि एल्गोरिथ्म हर पुनरावृत्ति के लिए गहराई से जाता है, आपको खुद को सीमाएं निर्धारित करनी होंगी। आमतौर पर स्टॉप इंडेक्स को एक प्रहरी द्वारा निरूपित किया जाता है। प्रहरी स्टॉप को बताने से आपको पता चल जाएगा कि कण किस स्तर पर रखा गया है। तो अधिकतम स्तर जिसे आप परिभाषित करना चाहते हैं, आपको प्रति आयाम कोशिकाओं की संख्या बताएगा। उदाहरण के लिए 6 पर अधिकतम स्तर के साथ आपके पास 2 ^ 6 = 64 है। आपके सिस्टम में (3D) 64x64x64 सेल होंगे। इसका मतलब यह भी है कि आपको पूर्णांक आधारित निर्देशांक का उपयोग करना होगा। यदि आप फ्लोट्स का उपयोग करते हैं तो आपको कन्वर्ट करना होगा coord.x = 64*float_x और इसी तरह।
यदि आप जानते हैं कि आपके पास आपके सिस्टम में कितनी कोशिकाएँ हैं तो आप अपनी सीमाओं को परिभाषित कर सकते हैं। क्या आप बाइनरी ऑक्ट्री का उपयोग करने की कोशिश कर रहे हैं?
चूंकि कण गति में हैं (उस CUDA उदाहरण में) आपको कोशिकाओं के बजाय कणों की संख्या पर समानांतर करने की कोशिश करनी चाहिए।
यदि आप निकटतम पड़ोसियों की सूची बनाना चाहते हैंआपको कणों को कोशिकाओं में मैप करना होगा। यह एक तालिका के माध्यम से किया जाता है जो कोशिकाओं द्वारा कणों को बाद में क्रमबद्ध किया जाता है। फिर भी आपको कणों के माध्यम से चलना चाहिए और इसके पड़ोसियों तक पहुंचना चाहिए।
आपके कोड के बारे में:
उस कोड के साथ समस्या यह है कि यह धीमा है। मुझे संदेह है कि नॉनलाइनियर जीपीयू मेमोरी एक्सेस (पार्टिक्युलरी पी [[[पी] .y] इनर-लूप के लिए सबसे अधिक है) में सुस्ती पैदा करता है।
डोनाल्ड नथ याद है। आपको मापना चाहिए कि बोतल गर्दन कहां है। आप NVCC प्रोफाइलर का उपयोग कर सकते हैं और टोंटी की तलाश कर सकते हैं। निश्चित नहीं है कि ओपनसीएल के पास प्रोफाइलर के रूप में क्या है।
// ugly boundary checking if ( dot(neigh_position<0, (float3)(1)) + dot(neigh_position>BOUNDARY, (float3)(1)) != 0) continue;
मुझे लगता है कि आपको इसे इस तरह से शाखा नहीं देना चाहिए, जब आप कॉल करते हैं तो शून्य वापस कैसे करें heavy_computation। यकीन नहीं हो रहा है, लेकिन हो सकता है कि आपके पास एक शाखा की भविष्यवाणी हो। किसी तरह उसे निकालने की कोशिश करें।
कोशिकाओं के समानांतर चलना एक अच्छा विचार हैकेवल अगर आपके पास कण डेटा तक कोई पहुंच नहीं है, अन्यथा आपको एटॉमिक्स का उपयोग करना होगा। यदि आप कोशिकाओं और पड़ोसियों तक पहुंच पढ़ते हैं, तो आप कण सीमा पर जाते हैं, लेकिन आप समानांतर में अपनी राशि बनाते हैं और आपको कुछ दौड़ संघनन प्रतिमान के लिए मजबूर नहीं किया जाता है।
इसके अलावा, कणों की संख्या / ग्रिड कोशिकाओं की संख्या का आदर्श अनुपात क्या है?
वास्तव में आपके एल्गोरिदम और पर निर्भर करता हैआपके डोमेन में कण पैकिंग, लेकिन आपके मामले में मैं कण व्यास के बराबर सेल आकार को परिभाषित करूंगा और बस आपको प्राप्त होने वाली कोशिकाओं की संख्या का उपयोग करूंगा।
इसलिए यदि आप Z- ऑर्डर का उपयोग करना चाहते हैं और अपनी पवित्र कब्र को प्राप्त करना चाहते हैं, तो पूर्णांक निर्देशांक का उपयोग करने का प्रयास करें और उन्हें हैश करें।
इसके अलावा बड़ी मात्रा में कणों का उपयोग करने की कोशिश करें।CUDA उदाहरणों के बारे में 65000 कणों का उपयोग करना चाहिए, आपको इस पर विचार करना चाहिए क्योंकि समानांतर तरीके ज्यादातर कुशल होते हैं; चल प्रसंस्करण इकाइयों का शोषण किया जाता है (कम बेकार धागे)।