सीयूडीए सरणी

Question

में संख्याओं के अवसरों की गणना करना मेरे पास सीयूडीए (आमतौर पर 1000000 तत्वों) के साथ GPU पर संग्रहीत हस्ताक्षरित पूर्णांक की एक सरणी है। मैं सरणी में प्रत्येक नंबर के अवसर को गिनना चाहता हूं। केवल कुछ विशिष्ट संख्याएं हैं (लगभग 10), लेकिन ये संख्याएं 1 से 1000000 तक फैली हो सकती हैं। 9/10 संख्याओं में से 0 हैं, मुझे उनकी गिनती की आवश्यकता नहीं है। परिणाम कुछ इस तरह दिखता है:

58458 -> 1000 occurences 
15 -> 412 occurences 

मैं एक कार्यान्वयन atomicAdd s का उपयोग करके है, लेकिन यह बहुत धीमी गति से (धागे का एक बहुत ही पते पर लिखें) है। क्या कोई तेज/कुशल विधि के बारे में जानता है?

Answer 1

आप संख्याओं को क्रमबद्ध करके पहले एक हिस्टोग्राम लागू कर सकते हैं, और फिर एक महत्वपूर्ण कमी कर सकते हैं।

सबसे सरल तरीका thrust::sort और फिर thrust::reduce_by_key का उपयोग करना होगा। यह परमाणुओं के आधार पर विज्ञापन कताई की तुलना में अक्सर तेज़ होता है। यहां एक example है।

Answer 2

मुझे लगता है कि आप सीयूडीए उदाहरणों, विशेष रूप से हिस्टोग्राम उदाहरणों में सहायता पा सकते हैं। वे जीपीयू कंप्यूटिंग एसडीके का हिस्सा हैं। आप इसे यहां देख सकते हैं http://developer.nvidia.com/cuda-cc-sdk-code-samples#histogram। उनके पास एल्गोरिदम समझाते हुए एक श्वेतपत्र भी है।

Answer 3

मैं दो तरीकों की तुलना कर रहा हूँ डुप्लिकेट प्रश्न thrust count occurence, अर्थात् में सुझाव दिया,

1. thrust::counting_iterator और thrust::upper_bound का उपयोग करना, हिस्टोग्राम थ्रस्ट उदाहरण का पालन करना;
2. thrust::unique_copy और thrust::upper_bound का उपयोग करना।

नीचे, कृपया एक पूरी तरह से काम किया उदाहरण खोजें।

#include <time.h>  // --- time 
#include <stdlib.h>  // --- srand, rand 
#include <iostream> 

#include <thrust\host_vector.h> 
#include <thrust\device_vector.h> 
#include <thrust\sort.h> 
#include <thrust\iterator\zip_iterator.h> 
#include <thrust\unique.h> 
#include <thrust/binary_search.h> 
#include <thrust\adjacent_difference.h> 

#include "Utilities.cuh" 
#include "TimingGPU.cuh" 

//#define VERBOSE 
#define NO_HISTOGRAM 

/********/ 
/* MAIN */ 
/********/ 
int main() { 

    const int N = 1048576; 
    //const int N = 20; 
    //const int N = 128; 

    TimingGPU timerGPU; 

    // --- Initialize random seed 
    srand(time(NULL)); 

    thrust::host_vector<int> h_code(N); 

    for (int k = 0; k < N; k++) { 
     // --- Generate random numbers between 0 and 9 
     h_code[k] = (rand() % 10); 
    } 

    thrust::device_vector<int> d_code(h_code); 
    //thrust::device_vector<unsigned int> d_counting(N); 

    thrust::sort(d_code.begin(), d_code.end()); 

    h_code = d_code; 

    timerGPU.StartCounter(); 

#ifdef NO_HISTOGRAM 
    // --- The number of d_cumsum bins is equal to the maximum value plus one 
    int num_bins = d_code.back() + 1; 

    thrust::device_vector<int> d_code_unique(num_bins); 
    thrust::unique_copy(d_code.begin(), d_code.end(), d_code_unique.begin()); 
    thrust::device_vector<int> d_counting(num_bins); 
    thrust::upper_bound(d_code.begin(), d_code.end(), d_code_unique.begin(), d_code_unique.end(), d_counting.begin()); 
#else 
    thrust::device_vector<int> d_cumsum; 

    // --- The number of d_cumsum bins is equal to the maximum value plus one 
    int num_bins = d_code.back() + 1; 

    // --- Resize d_cumsum storage 
    d_cumsum.resize(num_bins); 

    // --- Find the end of each bin of values - Cumulative d_cumsum 
    thrust::counting_iterator<int> search_begin(0); 
    thrust::upper_bound(d_code.begin(), d_code.end(), search_begin, search_begin + num_bins, d_cumsum.begin()); 

    // --- Compute the histogram by taking differences of the cumulative d_cumsum 
    //thrust::device_vector<int> d_counting(num_bins); 
    //thrust::adjacent_difference(d_cumsum.begin(), d_cumsum.end(), d_counting.begin()); 
#endif 

    printf("Timing GPU = %f\n", timerGPU.GetCounter()); 

#ifdef VERBOSE 
    thrust::host_vector<int> h_counting(d_counting); 
    printf("After\n"); 
    for (int k = 0; k < N; k++) printf("code = %i\n", h_code[k]); 
#ifndef NO_HISTOGRAM 
    thrust::host_vector<int> h_cumsum(d_cumsum); 
    printf("\nCounting\n"); 
    for (int k = 0; k < num_bins; k++) printf("element = %i; counting = %i; cumsum = %i\n", k, h_counting[k], h_cumsum[k]); 
#else 
    thrust::host_vector<int> h_code_unique(d_code_unique); 

    printf("\nCounting\n"); 
    for (int k = 0; k < N; k++) printf("element = %i; counting = %i\n", h_code_unique[k], h_counting[k]); 
#endif 
#endif 
} 

पहला दृष्टिकोण सबसे तेज़ दिखाया गया है।

First approach: 2.35ms 
First approach without thrust::adjacent_difference: 1.52 
Second approach: 4.67ms 

कृपया ध्यान दें, स्पष्ट रूप से आसन्न अंतर की गणना करने के लिए कोई सख्त जरूरत नहीं है कि, इस आपरेशन मैन्युअल हो सकता है के बाद से: एक NVIDIA GTX 960 कार्ड पर, मैं N = 1048576 सरणी तत्वों की एक संख्या के लिए निम्नलिखित समय पड़ा है यदि आवश्यक हो, तो कर्नेल प्रोसेसिंग के दौरान किया जाता है।