2013-05-19 104 views
3

我有一個給內核的輸入數組。每個線程都與數組中的一個值一起工作,並根據規則改變值或完全不改變它。在CUDA中模擬std :: bitset

我想很快就會發現輸入內存中是否有任何更改,如果有,我想很快找到發生此更改的位置(輸入數組的索引)。

我想過使用類似於位數組的東西。總的位數將等於線程總數。每個線程只會操縱一個位,所以最初這些位將被設置爲假,如果一個線程改變相應的輸入值,該位將變爲真。

爲了更清楚,讓我們假設我們有這個輸入數組稱爲A

1 9 3 9 4 5 

位的陣列將是以下

0 0 0 0 0 0 

因此,我們將有6個線程工作在輸入數組上。讓我們假設最終的輸入陣列將

1 9 3 9 2 5 

所以位的最後一個數組將是:

0 0 0 0 1 0 

我不想使用bool陣列,因爲每個值將取1字節的內存,這是相當多的,因爲我只想使用位來工作。

是否有可能實現這樣的目標?

我想創建char陣列,其中該陣列的每個值將有8位。但是,如果兩個線程想要更​​改數組第一個字符的不同位,該怎麼辦?即使位內部的變化位於不同位置,他們也必須以原子方式執行操作。因此,使用原子操作可能會破壞並行性,在這種情況下,不需要使用原子操作,它沒有任何意義,但必須使用,因爲使用chars數組的約束而不是更專門化像std::bitset

預先感謝您。

+0

這個問題看起來很像你:http://stackoverflow.com/questions/11042816/how-to-create-large-bit-array -in-CUDA – BenC

+0

感謝我讀的問題,但是回答這個問題不說我如何可以使用的比特的數組或東西,如果有像CUDA的'的std :: bitset'什麼。使用「bool」數組對我來說不是一個好主意,因爲我不能在GPU中使用太多的內存。 – ksm001

回答

3

我提供了一個遲到的回答這個問題從沒有答案的列表中刪除。

要做你想達到的目標,你可以定義一個長度爲N/32unsigned int的數組,其中N是你正在比較的數組的長度。然後就可以用atomicAdd寫這樣的陣列的每個比特,這取決於陣列的兩個元素是否相等或沒有。

在下面,我提供了一個簡單的例子:

#include <iostream> 

#include <thrust\device_vector.h> 

__device__ unsigned int __ballot_non_atom(int predicate) 
{ 
    if (predicate != 0) return (1 << (threadIdx.x % 32)); 
    else return 0; 
} 

__global__ void check_if_equal_elements(float* d_vec1_ptr, float* d_vec2_ptr, unsigned int* d_result, int Num_Warps_per_Block) 
{ 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    const unsigned int warp_num = threadIdx.x >> 5; 

    atomicAdd(&d_result[warp_num+blockIdx.x*Num_Warps_per_Block],__ballot_non_atom(!(d_vec1_ptr[tid] == d_vec2_ptr[tid]))); 
} 

// --- Credit to "C printing bits": http://stackoverflow.com/questions/9280654/c-printing-bits 
void printBits(unsigned int num){ 
    unsigned int size = sizeof(unsigned int); 
    unsigned int maxPow = 1<<(size*8-1); 
    int i=0; 
    for(;i<size;++i){ 
     for(;i<size*8;++i){ 
      // print last bit and shift left. 
      printf("%u ",num&maxPow ? 1 : 0); 
      num = num<<1; 
     }  
    } 
} 

void main(void) 
{ 
    const int N = 64; 

    thrust::device_vector<float> d_vec1(N,1.f); 
    thrust::device_vector<float> d_vec2(N,1.f); 

    d_vec2[3] = 3.f; 
    d_vec2[7] = 4.f; 

    unsigned int Num_Threads_per_Block  = 64; 
    unsigned int Num_Blocks_per_Grid  = 1; 
    unsigned int Num_Warps_per_Block  = Num_Threads_per_Block/32; 
    unsigned int Num_Warps_per_Grid   = (Num_Threads_per_Block*Num_Blocks_per_Grid)/32; 

    thrust::device_vector<unsigned int> d_result(Num_Warps_per_Grid,0); 

    check_if_equal_elements<<<Num_Blocks_per_Grid,Num_Threads_per_Block>>>((float*)thrust::raw_pointer_cast(d_vec1.data()), 
                      (float*)thrust::raw_pointer_cast(d_vec2.data()), 
                      (unsigned int*)thrust::raw_pointer_cast(d_result.data()), 
                      Num_Warps_per_Block); 

    unsigned int val = d_result[1]; 
    printBits(val); 
    val = d_result[0]; 
    printBits(val); 

    getchar(); 
}