在CUDA中模擬std :: bitset

Question

我有一個給內核的輸入數組。每個線程都與數組中的一個值一起工作，並根據規則改變值或完全不改變它。

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

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

爲了更清楚，讓我們假設我們有這個輸入數組稱爲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

預先感謝您。

Answer 1

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

要做你想達到的目標，你可以定義一個長度爲N/32的unsigned 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(); 
}