使用CUDA/Thrust使關鍵出現次數相等

Question

有沒有一種有效的方法來獲取排序後的鍵/值數組對並確保每個鍵使用CUDA Thrust庫都具有相同數量的元素？

例如，假設我們有下面的一對陣列：

ID: 1 2 2 3 3 3 
VN: 6 7 8 5 7 8 

如果我們想有兩個每個鍵的出現，這將是結果：

ID: 2 2 3 3 
VN: 7 8 5 7 

實際陣列將會更大，包含數百萬或更多的元素。我可以很容易地使用嵌套for循環來做到這一點，但我很想知道是否有更高效的方式來使用GPU轉換數組。 Thrust似乎可能很有用，但我沒有看到任何明顯的功能。

謝謝你的幫助！

Answer 1

警告：如果這是您計劃在GPU上執行的唯一操作，那麼我不會推薦它。從GPU複製數據的成本可能會超過使用GPU的任何可能的效率/性能優勢。

編輯：基於序列閾值可能會遠遠長於2的評論，我會建議一種替代方法（方法2），該方法應該比for-loop或brute-force更有效方法（方法1）。

一般來說，我會把這個問題放在一個名爲stream compaction的類別中。流壓縮通常是指採取一系列數據並將其減少爲較小的數據序列。

如果我們查看逆沖流壓實區域，可以對此問題進行處理的算法是thrust::copy_if()（特別是爲了方便起見，需要模板陣列的版本）。

方法1：

要想想並行這個問題，我們必須問自己在什麼條件下應該給定元素從輸入到輸出複製？如果我們能夠形式化這個邏輯，我們可以構造一個推力函數，我們可以傳遞給thrust::copy_if來指示它複製哪些元素。

對於給定的元件，對於序列長度= 2的情況下，我們可以構建完整的邏輯，如果我們知道：

1. 元件
2. 元件一個位置向右
3. 元件一個地方向左
4. 元素兩個地向左

基於以上，我們將需要拿出對於未定義上述第2,3或4項中的任何元素的「特例」邏輯。

忽略的特殊情況下，如果我們知道上面的4個項目，那麼我們可以構造必要的邏輯如下：

1. 如果元素到我的左邊是和我一樣，但元素2地方左邊是不同的，那麼我在輸出

2. 如果元素到我的左邊是比我不同的歸屬，但該元素在我右邊的是和我一樣，我在輸出屬於

3. 否則，我不屬於我n輸出

我會留給你來構造特殊情況下的必要邏輯。 （或者從我提供的代碼中反向工程）。

方法2：

對於長序列，方法1或方法1中的邏輯的一個for循環變體將產生至少1讀出的數據每序列長度的元件設置的。對於長序列（例如2000），這將是低效的。因此另一種可能的方法將是如下：

1. 向前生成在兩者exclusive_scan_by_key和反向方向上，使用ID值作爲密鑰，以及thrust::constant_iterator（值= 1）作爲用於掃描的值。對於給定的數據集，創建中間結果是這樣的：

   ID: 1 2 2 3 3 3 
   VN: 6 7 8 5 7 8 
   FS: 0 0 1 0 1 2 
   RS: 0 1 0 2 1 0 
   

其中FS和RS是向前的結果和反向掃描逐鍵。我們使用.rbegin()和.rend()reverse iterators生成反向掃描（RS）。請注意，爲了生成如上所述的RS序列，必須對反向掃描輸入和輸出執行此操作。

1. 我們的thrust::copy_if仿函數的邏輯變得相當簡單。對於給定元素，如果該元素的RS和FS值的總和大於或等於期望的最小序列長度（-1以考慮排他掃描操作）和，則FS值小於期望的最小序列長度，那麼該元素屬於輸出。

下面是兩種方法的整個例子，使用給定的數據，序列長度爲2：

$ cat t1095.cu 
#include <thrust/device_vector.h> 
#include <thrust/copy.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <iostream> 

#include <thrust/scan.h> 
#include <thrust/iterator/constant_iterator.h> 

struct copy_func 
{ 
    int *d; 
    int dsize, r, l, m, l2; 
    copy_func(int *_d, int _dsize) : d(_d),dsize(_dsize) {}; 
    __host__ __device__ 
    bool operator()(int idx) 
    { 
    m = d[idx]; 
    // handle typical case 
    // this logic could be replaced by a for-loop for sequences of arbitrary length 
    if ((idx > 1) && (idx < dsize-1)){ 
     r = d[idx+1]; 
     l = d[idx-1]; 
     l2 = d[idx-2]; 
     if ((r == m) && (m != l)) return true; 
     if ((l == m) && (m != l2)) return true; 
     return false;} 
    // handle special cases 
    if (idx == 0){ 
     r = d[idx+1]; 
     return (r == m);} 
    if (idx == 1){ 
     r = d[idx+1]; 
     l = d[idx-1]; 
     if (l == m) return true; 
     else if (r == m) return true; 
     return false;} 
    if (idx == dsize-1){ 
     l = d[idx-1]; 
     l2 = d[idx-2]; 
     if ((m == l) && (m != l2)) return true; 
     return false;} 
    // could put assert(0) here, should never get here 
    return false; 
    } 
}; 

struct copy_func2 
{ 
    int thresh; 
    copy_func2(int _thresh) : thresh(_thresh) {}; 
    template <typename T> 
    __host__ __device__ 
    bool operator()(T t){ 
    return (((thrust::get<0>(t) + thrust::get<1>(t))>=(thresh-1)) && (thrust::get<0>(t) < thresh)); 
    } 
}; 

int main(){ 

    const int length_threshold = 2; 
    int ID[] = {1,2,2,3,3,3}; 
    int VN[] = {6,7,8,5,7,8}; 
    int dsize = sizeof(ID)/sizeof(int); 
    // we assume dsize > 3 
    thrust::device_vector<int> id(ID, ID+dsize); 
    thrust::device_vector<int> vn(VN, VN+dsize); 

    thrust::device_vector<int> res_id(dsize); 
    thrust::device_vector<int> res_vn(dsize); 
    thrust::counting_iterator<int> idx(0); 

    //method 1: sequence length threshold of 2 

    int rsize = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(id.begin(), vn.begin())), thrust::make_zip_iterator(thrust::make_tuple(id.end(), vn.end())), idx, thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())), copy_func(thrust::raw_pointer_cast(id.data()), dsize)) - thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())); 

    std::cout << "ID: "; 
    thrust::copy_n(res_id.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl << "VN: "; 
    thrust::copy_n(res_vn.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 

    //method 2: for arbitrary sequence length threshold 
    thrust::device_vector<int> res_fs(dsize); 
    thrust::device_vector<int> res_rs(dsize); 
    thrust::exclusive_scan_by_key(id.begin(), id.end(), thrust::constant_iterator<int>(1), res_fs.begin()); 
    thrust::exclusive_scan_by_key(id.rbegin(), id.rend(), thrust::constant_iterator<int>(1), res_rs.begin()); 
    rsize = thrust::copy_if(thrust::make_zip_iterator(thrust::make_tuple(id.begin(), vn.begin())), thrust::make_zip_iterator(thrust::make_tuple(id.end(), vn.end())), thrust::make_zip_iterator(thrust::make_tuple(res_fs.begin(), res_rs.rbegin())), thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())), copy_func2(length_threshold)) - thrust::make_zip_iterator(thrust::make_tuple(res_id.begin(), res_vn.begin())); 

    std::cout << "ID: "; 
    thrust::copy_n(res_id.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl << "VN: "; 
    thrust::copy_n(res_vn.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

$ nvcc -o t1095 t1095.cu 
$ ./t1095 
ID: 2 2 3 3 
VN: 7 8 5 7 
ID: 2 2 3 3 
VN: 7 8 5 7 

注：

1. 的copy_func實現測試邏輯對於給定的元素。它通過函數初始化參數接收該元素的索引（通過模板）以及指向設備上的數據的指針以及數據的大小。變量r,m,l和l2分別指向右側的元素，我自己，左側的元素和左側的兩個元素。

2. 我們將一個指針傳遞給函數的ID數據。這允許仿函數檢索測試邏輯的（最多）4個必要元素。這避免了一個棘手的建設推拉:: zip_iterator提供所有這些值。請注意，仿函數中這些元素的讀取應該很好地融合在一起，因此效率相當高，並且也受益於緩存。

3. 我不認爲這是無缺陷的。我認爲我的測試邏輯是正確的，但有可能我沒有。至少，您應該驗證該部分代碼的邏輯正確性。我的目的不是要給你一個黑盒子的代碼，而是要演示如何思考你通過問題的方式。

4. 這種方法對於長於2的關鍵序列可能會很麻煩。在這種情況下，我會建議方法2.（如果您已經有了一個順序的for循環來實現必要的邏輯，那麼您可以放棄一個修改這種for-loop應該可能仍然受益於來自緩存的合併訪問和相鄰訪問）。