CUDA पर O (N) में एक सरणी को क्रमबद्ध करें

परिचय

किसी भी तरह से कम से कम कोड और उच्चतम संभव गति के साथ एक GPU का उपयोग करके तार की एक अनूठी सरणी को सॉर्ट करने का कार्य था ...

इस पोस्ट में मैं इसके समाधान के मुख्य विचार का वर्णन करूंगा। इस पोस्ट में छँटाई सरणी के तत्व संख्याएँ हैं।

एक छोटे से सरणी के अनूठे तत्वों के साथ मामला

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

CUDA सॉर्टिंग एल्गोरिदम के लिए एक नेटवर्क खोज के विभिन्न परिणाम मिले। यहाँ सबसे दिलचस्प है । एक ड्राइंग है



, जो दर्शाता है कि QSORT एल्गोरिथ्म द्वारा सबसे अच्छा परिणाम दिया गया था, जो O (NlogN) से O (N ^ 2) तक ऑर्डर जटिलता देता है। यद्यपि GPU पर समानांतरण ने लेख में सबसे अच्छा परिणाम प्राप्त किया, लेकिन इस प्रश्न में क्रेप इस वीडियो के लिए वीडियो कार्ड के संसाधनों का उपयोग करने का सबसे अच्छा तरीका नहीं है (यह विशेष रूप से दिए गए कोड के आकार के लिए डरावना था)। निम्नलिखित समस्या के समाधान का वर्णन करता है, अनिवार्य रूप से सबसे खराब स्थिति में ओ (एन) की समय जटिलता की जटिलता के साथ "एक पंक्ति"।





तो, हमारे पास क्या है: आधुनिक वीडियो कार्ड में 500 * 32 से अधिक धागे (GTX275) हैं, एक साथ काम कर रहे हैं, लगभग 15,000।

सॉर्टिंग टास्क के लिए स्थितियां इस प्रकार थीं: 1000 तत्वों तक की सरणी (इसके बाद, एक बड़े मूल्य के साथ सामान्य मामला माना जाता है, लेकिन मेरे पास था)। ओ (NlogN) समय के साथ इस तरह के एक सरणी को छांटना, 15,000 धागे होने से किसी भी तरह बेकार लग रहा था।



पेपिरस तर्क के परिणामस्वरूप, यह पाया गया कि, सबसे अच्छी गति प्राप्त करने के लिए, और यह ओ (एन) है, यह आवश्यक है कि एक साथ काम करने वाला प्रत्येक धागा ओ (एन) संचालन से अधिक नहीं करे। यानी एक बार हमारे एरे से होकर गुजरे। और पास का परिणाम एक कुंजी होना चाहिए जो परिणामी सॉर्ट किए गए सरणी में एक तत्व की स्थिति निर्धारित करता है।

स्वयं थ्रेड नंबर tid (आमतौर पर 0..MaxThreads-1 से) को जानने के बाद, परिणाम के रूप में tid तत्व की स्थिति का पता लगाने का विचार किया गया था।



मूल सरणी को A [] = {1,5,2,4,7} के रूप में लेते हुए, हम परिणामस्वरूप परिणाम B [] में तत्व A [tid] की स्थिति का पता लगाने के लिए प्रत्येक थ्रेड के लिए कार्य निर्धारित करते हैं।



ऐसा करने के लिए, हम एक सहायक सरणी K [] शुरू करते हैं - परिणामस्वरूप तत्व A [tid] की स्थिति। ArraySize - ए [] में तत्वों की संख्या।

अब स्रोत कोड स्वयं CUDA C (.cu फ़ाइल) में प्रस्तावित एल्गोरिथ्म के कार्यान्वयन का एक टुकड़ा है:

1. MainKernel ( kernel.cu) ArraySize .

MainKernel <<<grids, threads>>> (A,B,ArraySize,MaxThreads);



2. CUDA- MainKernel ( kernel.cu) A[], ArraySize, MaxThreads, B[] K[].

unsigned long tid = blockIdx.x*blockDim.x + threadIdx.x; //

for (unsigned long i = 0;i<ArraySize;++i) // O(N)

if (A[i]<A[tid])

++K[tid]; // O(1)



3. __syncthreads(); // , .. tid-



4. B[K[tid]] = A[tid]; // ( O(1), 3 qsort, swap ..)

* , K[]={0,0,2,3,4}.

.4 , ( O(1)). .







और वह सब, परिणामी सरणी B [] O (N) द्वारा सॉर्ट किया जाता है, बशर्ते कि उपलब्ध थ्रेड्स की संख्या N के बराबर या उससे अधिक हो।



एक बड़े सरणी के अनूठे तत्वों के साथ मामला

अब, यदि आप सरणी के आयाम को बढ़ाते हैं, तो आप उसी दृष्टिकोण का उपयोग कर सकते हैं, केवल एक धागा पदों के लिए दिखेगा

[ArraySize / MaxThreads] - (तत्वों का पूर्णांक भाग)। इस स्थिति में, एल्गोरिथ्म का रनिंग टाइम बदल जाएगा [ArraySize / MaxThreads] * O (N), यानी यह कई बार बन जाएगा क्योंकि सरणी में कई तत्व उपलब्ध थ्रेड्स की संख्या से अधिक हैं।



इस स्थिति के लिए, एल्गोरिथ्म के कार्यान्वयन के लिए फार्म होगा:

1. MainKernel ( kernel.cu) [ArraySize/MaxThreads] .

MainKernel <<<grids, threads>>> (A,B,ArraySize,MaxThreads);



2. CUDA- MainKernel ( kernel.cu) A[], ArraySize, MaxThreads, B[] K[].

unsigned long tid = blockIdx.x*blockDim.x + threadIdx.x; //



2.5. , .

unsigned long c = ArraySize/MaxThreads;

for (unsigned long j = 0;j<c;++j) {

unsigned int ttid = c*tid + j; // -



for (unsigned long i = 0;i<ArraySize;++i)

if (A[i]<A[ttid])

++K[ttid]; //



3. __syncthreads(); // , .. tid-



4. B[K[ttid]] = A[ttid]; //

} //







निष्कर्ष 1

एक पंक्ति में लिखी संख्या tid के साथ एक स्ट्रीम के लिए अंतिम छँटाई एल्गोरिथ्म का वादा किया:

for (unsigned long i = 0;i<ArraySize;++i) if (A[i]<A[tid]) ++K[tid];





यह था: ए [i] - एक अनसुलझा सरणी। अब K [tid] - परिणामस्वरूप स्रोत तत्वों की स्थिति सूचकांकों में शामिल है। B [K [tid]] = A [tid] और B [ArraySize-K [tid] -1] = A [tid] के बीच चयन करके छँटाई क्रम को बदलना आसान है।



शुरुआत से डेटा के साथ उदाहरण

यह था: A [] = {1,5,2,4,7}। ArraySize = 5. हम 5 प्रवाह शुरू करते हैं।

स्ट्रीम 0. tid = 0. सेट K [tid] = K [0] A [tid] = A [0] = 1: K [0] = 0 के लिए

स्ट्रीम 1. tid = 1. सेट K [tid] = K [1] A [tid] = A [1] = 5: K [1] = 3

धारा 2. tid = 2. A [2] = 2। के [२] = १

स्ट्रीम 3. tid = 3. A [3] = 4। के [३] = २

स्ट्रीम 4. tid = 4. A [4] = 7। के [४] = ४

वह सब है। अब, चूंकि सभी धागे एक ही समय में चल रहे थे, हमने केवल सरणी - ओ (एन) के माध्यम से छंटनी में समय बिताया। और उन्हें K [] = {0,3,1,2,4} मिला।



विशेषताएं

जब [ArraySize / MaxThreads] काफी बड़ा हो गया है > ArraySize, यह एल्गोरिथ्म O (N ^ 2) से अधिक जटिलता लेता है।

मैं इसका उपयोग करने की अनुशंसा नहीं करता हूं जब यह पहले से ज्ञात नहीं है कि मूल सरणी के तत्व अद्वितीय हैं या नहीं।



निष्कर्ष २

इस प्रकार, मूल समस्या ओ (एन) के लिए सबसे खराब और ओ (एन) सर्वोत्तम मामलों में हल की गई थी (त्वरण केवल पाठ्यक्रम के एन प्रवाह के कारण था)। और फिर बशर्ते कि थ्रेड की उपलब्ध संख्या सॉर्ट किए गए सरणी के तत्वों की संख्या से कम नहीं है।



इस पोस्ट को नौसिखिए एल्गोरिथ्म शोधकर्ताओं को GPGPU पर तेजी से कंप्यूटिंग के लिए दिलचस्प वैचारिक समाधान लिखने में मदद करनी चाहिए। साहित्य से, मैं कॉर्मेन की एल्गोरिदम की पुस्तक और दिलचस्प समस्याओं के पोर्टल को नोट करता हूं projecteuler.net।



UPD1। (12/17/2010 1:40 PM)

इस प्रक्रिया में, यह पता चला कि यदि मूल सरणी A [] में 16KB मेमोरी होती है, तो आपको अभी भी स्टोर करना होगा और परिणाम B [] __sared__ मेमोरी में काम नहीं करेगा (और हम इसका उपयोग CUDA में एक स्तर 2 कैश के रूप में गणना करने में तेजी लाने के लिए करते हैं) - क्योंकि __sared__ का आकार केवल 16KB है (GTX275 के लिए)।

इसलिए, एल्गोरिथ्म को "समानांतर स्वैप" के साथ पूरक किया जाता है, जहां परिणाम मूल ए [] होता है:

1. for (unsigned long i = 0;i<ArraySize;i++)

if (A[i]<A[tid])

K[tid]++; //



2. __syncthreads(); // !



3. unsigned long B_tid = A[tid]; //, ""



4. __syncthreads(); // ""



5. A[K[tid]] = B_tid; //





अब मूल सरणी A [] कम मेमोरी का उपयोग करके सॉर्ट किया गया है।



UPD2। (12/17/2010 16:34, विषय )

__Sared__ मेमोरी और छोटे अनुकूलन के लिए कम पहुँच के कारण तेज़ विकल्प।

0. unsigned int K_tid = 0;



1. unsigned long B_tid = A[tid]; //, ""



2. for (unsigned long i = 0;i<ArraySize;++i)

if (A[i]<B_tid)

++K_tid; //



3. __syncthreads(); // ""



4. A[K_tid] = B_tid; //







UPD3 । ( 12.20.2010 14:16) मैं उदाहरण के लिए स्रोत कोड प्रदान करता हूं: main.cpp , kernel.cu (सरणी का आकार वीडियो कार्ड की __sared__ मेमोरी के भीतर होना चाहिए)।



UPD4 । (12/23/2010 2:56 अपराह्न) जटिलता O (N) के संबंध में: प्रस्तावित एल्गोरिथ्म रैखिक समय जटिलता Θ (n + k) © विकी के साथ गणना छंटाई एल्गोरिथ्म का एक प्रकार है।