使用 GPU 获取数字列表的统计信息

Question

我在一个文件上有几个数字列表。例如，

.333, .324, .123 , .543, .00054
.2243, .333, .53343 , .4434

现在，我想使用 GPU 获取每个数字出现的次数。我相信这在 GPU 上会比在 CPU 上更快，因为每个线程都可以处理一个列表。我应该在 GPU 上使用什么数据结构来轻松获得上述计数。例如，对于上述情况，答案将如下所示：

.333 = 2 times in entire file
.324 = 1 time

等等。

我正在寻找一个通用的解决方案。不是仅适用于具有特定计算能力的设备

只需编写 Pavan 建议的内核，看看我是否有效地实现了它：

int uniqueEle = newend.valiter – d_A;

int* count;
cudaMalloc((void**)&count, uniqueEle * sizeof(int)); // stores the count of each unique element
int TPB = 256;
int blocks = uniqueEle + TPB -1 / TPB;
//Cast d_I to raw pointer called d_rawI
launch<<<blocks,TPB>>>(d_rawI,count,uniqueEle);

__global__ void launch(int *i, int* count, int n){
    int id = blockDim.x * blockIdx.x + threadIdx.x;
    __shared__ int indexes[256];
    if(id < n ){
        indexes[threadIdx.x] = i[id];
        //as occurs between two blocks
        if(id % 255 == 0){
            count[indexes] = i[id+1] - i[id];
        }
    }
    __syncthreads();
    if(id < ele - 1){
        if(threadIdx.x < 255)
            count[id] = indexes[threadIdx.x+1] – indexes[threadIdx.x];

    }
}

问题：如何修改此内核以使其处理任意大小的数组。即处理线程总数时的情况

Answer 1

这是我在matlab中的代码

A = [333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434];
[values, locations] = unique(A);   % Find unique values and their locations
counts = diff([0, locations]);     % Find the count based on their locations

在纯 cuda 中没有简单的方法可以做到这一点，但您可以使用现有的库来做到这一点。

1)Thrust

它还随 CUDA 4.0 的 CUDA 工具包一起提供。

matlab代码可以通过以下函数粗略翻译成推力。我对推力不太精通，但我只是想让您了解一下要查看的例程。

float _A[] = {.333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434};
int _I[] = {0, 1, 2, 3, 4, 5, 6, 7, 8};
float *A, *I; 
// Allocate memory on device and cudaMempCpy values from _A to A and _I to I
int num = 9;
// Values vector
thrust::device_vector<float>d_A(A, A+num);
// Need to sort to get same values together    
thrust::stable_sort(d_A, d_A+num);
// Vector containing 0 to num-1
thrust::device_vector<int>d_I(I, I+num);
// Find unique values and elements
thrust::device_vector<float>d_Values(num), d_Locations(num), d_counts(num);
// Find unique elements
thrust::device_vector<float>::iterator valiter;
thrust::device_vector<int>::iterator idxiter;
thrust::pair<valiter, idxiter> new_end;
new_end = thrust::unique_by_key(d_A, d_A+num, d_I, d_Values, d_Locations);

您现在有了每个唯一值的第一个实例的位置。您现在可以启动内核来查找 d_Locations 中从 0 到 new_end 的相邻元素之间的差异。从 num 中减去最终值以获得最终位置的计数。

编辑（添加通过聊天提供的代码）

这里是差异代码需要如何完成

#define MAX_BLOCKS 65535
#define roundup(A, B) = (((A) + (B) - 1) / (B))

int uniqueEle = newend.valiter – d_A;
int* count;
cudaMalloc((void**)&count, uniqueEle * sizeof(int));

int TPB = 256;
int num_blocks = roundup(uniqueEle, TPB);
int blocks_y = roundup(num_blocks, MAX_BLOCKS);
int blocks_x = roundup(num_blocks, blocks_y);
dim3 blocks(blocks_x, blocks_y);

kernel<<<blocks,TPB>>>(d_rawI, count, uniqueEle);

__global__ void kernel(float *i, int* count, int n)
{
int tx = threadIdx.x;
int bid = blockIdx.y * gridDim.x + blockIdx.x;
int id = blockDim.x * bid + tx;
__shared__ int indexes[256];

if (id < n) indexes[tx] = i[id];
__syncthreads();

if (id < n - 1) {
if (tx < 255) count[id] = indexes[tx + 1] - indexes[tx];
else count[id] = i[id + 1] - indexes[tx];
}

if (id == n - 1) count[id] = n - indexes[tx];
return;
}

2) ArrayFire

这是一个易于使用、免费基于数组的库。

您可以在 ArrayFire 中执行以下操作。

using namespace af;
float h_A[] = {.333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434};
int num = 9;
// Transfer data to device
array A(9, 1, h_A);
array values, locations, original;
// Find the unique values and locations
setunique(values, locations, original, A);
// Locations are 0 based, add 1.
// Add *num* at the end to find count of last value. 
array counts = diff1(join(locations + 1, num));

披露：我为开发此软件的 AccelerEyes 工作。

Answer 2

这里有一个解决方案：

__global__ void counter(float* a, int* b, int N)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    if(idx < N)
    {
        float my = a[idx];
        int count = 0;
        for(int i=0; i < N; i++)
        {
            if(my == a[i])
                count++;
        }

        b[idx]=count;
    }
}

int main()
{

    int threads = 9;
    int blocks = 1;
    int N = blocks*threads;
    float* h_a;
    int* h_b;
    float* d_a;
    int* d_b;

    h_a = (float*)malloc(N*sizeof(float));
    h_b = (int*)malloc(N*sizeof(int));

    cudaMalloc((void**)&d_a,N*sizeof(float));
    cudaMalloc((void**)&d_b,N*sizeof(int));

    h_a[0]= .333f; 
    h_a[1]= .324f;
    h_a[2]= .123f;
    h_a[3]= .543f;
    h_a[4]= .00054f;
    h_a[5]= .2243f;
    h_a[6]= .333f;
    h_a[7]= .53343f;
    h_a[8]= .4434f;

    cudaMemcpy(d_a,h_a,N*sizeof(float),cudaMemcpyHostToDevice);

    counter<<<blocks,threads>>>(d_a,d_b,N);

    cudaMemcpy(h_b,d_b,N*sizeof(int),cudaMemcpyDeviceToHost);

    for(int i=0; i < N; i++)
    {
        printf("%f = %d times\n",h_a[i],h_b[i]);
    }

    cudaFree(d_a);
    cudaFree(d_b);
    free(h_a);
    free(h_b);
    getchar();
    return 0;
}

Answer 3

要回答这个问题的最新附录 - 将完成 thrust 方法 proposed by Pavan 的差异内核可能看起来像这样：

template<int blcksz>
__global__ void diffkernel(const int *i, int* count, const int n) { 
    int id = blockDim.x * blockIdx.x + threadIdx.x; 
    int strd = blockDim.x * gridDim.x;
    int nmax = blcksz * ((n/blcksz) + ((n%blcksz>0) ? 1 : 0));

    __shared__ int indices[blcksz+1]; 

    for(; id<nmax; id+=strd) {
        // Data load
        indices[threadIdx.x] = (id < n) ? i[id] : n; 
        if (threadIdx.x == (blcksz-1)) 
            indices[blcksz] = ((id+1) < n) ? i[id+1] : n; 

        __syncthreads(); 

        // Differencing calculation
        int diff = indices[threadIdx.x+1] - indices[threadIdx.x];

        // Store
        if (id < n) count[id] = diff;

        __syncthreads(); 
    }
}