CUDA에서 std :: bitset 에뮬레이트

Question

커널에 입력 배열이 있습니다. 각 스레드는 배열의 한 값으로 작업하고 규칙에 따라 값을 변경하거나 전혀 변경하지 않습니다.

입력 메모리 내부에 변경 사항이있는 경우 매우 빨리 찾고 싶습니다.이 경우 변경 사항 (입력 배열의 인덱스)이 발생한 위치를 매우 빠르게 찾고 싶습니다.

비트 배열과 같은 것을 사용하려고 생각했습니다. 총 비트 수는 전체 스레드 수와 같습니다. 각 스레드는 하나의 비트 만 조작하므로 처음에는 비트가 false로 설정됩니다. 스레드가 해당 입력 값을 변경하면 비트는 참이됩니다.

의 우리가이 입력 배열이 비트의 배열이 작업은 다음

0 0 0 0 0 0 

그래서 우리가 가지고있는 것 (6 개) 스레드 것 A

1 9 3 9 4 5 

라는 있다고 가정하자, 좀 더 명확하게하기 위해서 입력 배열에. 의 최종 입력 배열이 될 것이라고 가정하자

1 9 3 9 2 5 

그래서 비트의 최종 배열은 다음과 같습니다

0 0 0 0 1 0 

나는 각각의 값이 걸릴 것이기 때문에 bool의 배열을 사용하지 않으 내가 비트만을 사용하기 때문에 꽤 많은 메모리의 1 바이트.

이렇게 할 수 있습니까?

배열의 각 값이 8 비트가되는 char 배열을 만들려고했습니다. 그러나 두 개의 스레드가 배열의 첫 문자의 다른 비트를 변경하려면 어떻게해야합니까? 비트 내의 변경이 다른 위치에있을지라도 그들은 원자 적으로 연산을 수행해야 할 것입니다. 따라서 원자 연산을 사용하면 병렬 처리가 중단 될 수 있습니다.이 경우 원자 연산을 사용하는 것은 필요하지 않지만 이해가되지 않지만보다 전문화 된 것 대신 char 배열을 사용하는 제약 때문에 사용해야합니다. like std::bitset

감사합니다.

Answer 1

응답하지 않은 목록에서이 질문을 삭제하려면이 질문에 대한 늦은 답변을 제공하고 있습니다.

달성하고자하는 것을 수행하려면 인 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(); 
}