2011-04-19 331 views
4

我有幾個塊,每個塊在整數數組的單獨部分執行。舉個例子:從數組[0]到數組[9]阻塞一個,從數組[10]到數組[20]阻塞兩個。CUDA:獲取數組中的最大值及其索引

什麼是我可以在陣列的每個塊的最大值的指數的最佳途徑?

實施例塊之一的[0]到[10]具有下列值:
5 10 2 3 4 34 56 3 9 10

所以56是在索引6.

我不能使用共享存儲器,因爲最大的值數組的大小可能非常大。因此它不適合。有沒有任何圖書館可以讓我這麼快?

我知道的簡化算法,但我認爲我的情況是不同的,因爲我想獲得最大的元素的索引。

+1

只是爲了理解。你在數組中有56個,你說34是最大的值。這是一個錯字嗎? – dubnde 2011-04-19 17:42:17

+0

你忘了提及你正在使用'CUDA'設置。 – 2011-04-19 18:39:09

回答

2

如果我的理解正是你想要的是:獲取裏面的最大值的數組A指數。

如果這是真的話,我會建議你使用推力庫:

這裏是你會怎麼做:

#include <thrust/device_vector.h> 
#include <thrust/tuple.h> 
#include <thrust/reduce.h> 
#include <thrust/fill.h> 
#include <thrust/generate.h> 
#include <thrust/sort.h> 
#include <thrust/sequence.h> 
#include <thrust/copy.h> 
#include <cstdlib> 
#include <time.h> 

using namespace thrust; 

// return the biggest of two tuples 
template <class T> 
struct bigger_tuple { 
    __device__ __host__ 
    tuple<T,int> operator()(const tuple<T,int> &a, const tuple<T,int> &b) 
    { 
     if (a > b) return a; 
     else return b; 
    } 

}; 

template <class T> 
int max_index(device_vector<T>& vec) { 

    // create implicit index sequence [0, 1, 2, ...) 
    counting_iterator<int> begin(0); counting_iterator<int> end(vec.size()); 
    tuple<T,int> init(vec[0],0); 
    tuple<T,int> smallest; 

    smallest = reduce(make_zip_iterator(make_tuple(vec.begin(), begin)), make_zip_iterator(make_tuple(vec.end(), end)), 
         init, bigger_tuple<T>()); 
    return get<1>(smallest); 
} 

int main(){ 

    thrust::host_vector<int> h_vec(1024); 
    thrust::sequence(h_vec.begin(), h_vec.end()); // values = indices 

    // transfer data to the device 
    thrust::device_vector<int> d_vec = h_vec; 

    int index = max_index(d_vec); 

    std::cout << "Max index is:" << index <<std::endl; 
    std::cout << "Value is: " << h_vec[index] <<std::endl; 

    return 0; 
} 
+0

我想她問她是否可以打電話給max_index(d_vec);從內核裏面?在設備上? – scatman 2011-04-20 05:40:09

0

除了建議使用推力,你也可以使用CUBLAS cublasIsamax函數。

0

相比於共享存儲器的陣列的大小几乎是無關緊要的,因爲線程的每個塊中的數是限制因素,而不是在陣列的大小。一種解決方案是讓每個線程塊的大小與線程塊的大小相同。也就是說,如果你有512個線程,那麼塊n將看着數組[n]到數組[n + 511]。每個塊都會減少以找到該部分中最高的成員。然後,將每個部分的最大值返回給主機,並執行簡單的線性搜索以找到整個陣列中的最高值。每次減少GPU不會將線性搜索減少512倍。根據陣列的大小,您可能希望在將數據恢復前進行更多減少。 (如果您的陣列尺寸爲3 * 512^10,則可能需要對GPU執行10次減少操作,並讓主機搜索其餘3個數據點。)

0

有一點需要注意,最大值加索引減少是因爲如果數組中存在多個相同值的最大元素,即在您的示例中,如果有2個或更多值等於56,那麼返回的索引將不是唯一的並且可能不同在代碼的每次運行中,因爲GPU上的線程排序的時序不是確定性的。

爲了解決這樣的問題,你可以使用一個唯一的順序設置指標,如線程ID + threadsperblock *塊標識,否則元素的索引位置,如果這是唯一的。然後最大考驗是沿着這些線路:

if(a>max_so_far || a==max_so_far && order_a>order_max_so_far) 
{ 
    max_so_far = a; 
    index_max_so_far = index_a; 
    order_max_so_far = order_a; 
} 

(索引和順序可以是相同的變量,取決於應用程序。)

2

這將不利於原始的海報,但對於那些誰來到這個頁面尋找答案我會建議使用推力已經有一個功能thrust :: max_element完全是 - 返回最大元素的索引。還提供了min_element和minmax_element函數。詳情請參閱推力文檔here

相關問題