4

カーネルに与えられる入力配列があります。各スレッドは配列の 1 つの値で動作し、ルールに従って値を変更するか、まったく変更しません。

入力メモリ内に変更があった場合は、後で非常に迅速に確認したいと思います。変更があった場合は、この変更が発生した場所 (入力配列のインデックス) を非常に迅速に確認したいと考えています。

ビットの配列のようなものを使用することを考えました。ビットの合計量は、スレッドの合計量に等しくなります。各スレッドは 1 ビットのみを操作するため、最初はビットが false に設定されます。スレッドが対応する入力値を変更すると、ビットは true になります。

より明確にするために、この入力配列が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ビットになる配列を作成することを考えました。しかし、2 つのスレッドが配列の最初の文字の異なるビットを変更したい場合はどうなるでしょうか? ビット内の変更が異なる場所にある場合でも、操作をアトミックに実行する必要があります。したがって、アトミック操作を使用すると、おそらく並列処理が中断されます。この場合、アトミック操作を使用する必要はありません。意味がありませんが、より特殊なものではなく文字の配列を使用するという制約があるため、使用する必要があります。のようにstd::bitset

前もって感謝します。

4

1 に答える 1

3

未回答のリストから削除するために、この質問に対する遅い回答を提供しています。

達成したいことを行うには、 length のunsigned ints の配列を定義できますN/32。ここNで、 は比較する配列の長さです。次にatomicAdd、配列の 2 つの要素が等しいかどうかに応じて、そのような配列の各ビットを書き込むために を使用できます。

以下に簡単な例を示します。

#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();
}
于 2014-05-26T20:48:07.580 に答える