时间戳

首页 > 解决方案 > 高性能计算和线程标识符的保存

c++ - 高性能计算和线程标识符的保存

问题描述

我编写网格步长循环来进行高性能计算，其中 N 很大，例如 long long N 1<<36，甚至更多。从总网格中我只需要一些索引，它们必须满足定义条件。

__global__ void Indexes(int *array, int N) {
int  index  = blockIdx.x * blockDim.x + threadIdx.x;
while( index<N)
    {
       if (condition)
       {....//do something to save index in array}  
    index += blockDim.x * gridDim.x;            
    }
}

当然，也可以使用 Thrust，它允许同时拥有主机和设备阵列。但在这种情况下，计算显然是非常无效的，因为需要先创建很多不需要的元素，然后再删除这些元素。

将索引直接保存在设备中的数组中以传入CPU的最有效方法是什么？

标签： c++cudahpc

解决方案

如果您的输出相对密集（即很多索引和相对较少的零），那么注释中建议的流压缩方法是一个很好的解决方案。有很多现成的流压缩实现，您可能可以根据自己的目的进行调整。

如果您的输出是稀疏的，因此您需要为大量输入保存相对较少的索引，那么流压缩不是一个很好的解决方案，因为它会浪费大量的 GPU 内存。在这种情况下（您可以粗略估计输出索引数量的上限），如下所示：

template <typename T>
struct Array 
{
    T*  p;
    int Nmax;
    int* next;  

    Array() = default;

    __host__ __device__ 
    Array(T* _p, int _Nmax, int* _next) : p(_p), Nmax(_Nmax), next(_next) {};

    __device__
    int append(T& val)
    {
        int pos = atomicAdd(next, 1);
        if (pos > Nmax) {
            atomicExch(next, Nmax);
            return -1;
        } else {           
            p[pos] = val;
            return pos;
        }
    };
};

可能更合适。在这里，想法是在输出数组中使用原子递增的位置来跟踪线程应该存储其索引的位置。如果您填充了索引数组，代码将发出信号，并且您可以从中制定重新启动策略以停止当前内核，然后从您能够存储的最后一个已知索引开始。

一个完整的例子：

$ cat append.cu 

#include <iostream>
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/copy.h>

namespace AppendArray
{
    template <typename T>
    struct Array 
    {
        T*  p;
        int Nmax;
        int* next;  

        Array() = default;

        __host__ __device__ 
        Array(T* _p, int _Nmax, int* _next) : p(_p), Nmax(_Nmax), next(_next) {};

        __device__
        int append(T& val)
        {
            int pos = atomicAdd(next, 1);
            if (pos > Nmax) {
                atomicExch(next, Nmax);
                return -1;
            } else {           
                p[pos] = val;
                return pos;
            }
        };
    };
}

    __global__ 
void kernelfind(int* input, int N, AppendArray::Array<int> indices)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    for(; idx < N; idx += gridDim.x*blockDim.x) {
        if (input[idx] % 10000 == 0) {
            if (indices.append(idx) < 0) return;
        }
    }
}

int main()
{
    const int Ninputs =  1 << 20;
    thrust::device_vector<int> inputs(Ninputs);
    thrust::counting_iterator<int> vals(1);
    thrust::copy(vals, vals + Ninputs, inputs.begin());
    int* d_input = thrust::raw_pointer_cast(inputs.data());

    int Nindices =  Ninputs >> 12;
    thrust::device_vector<int> indices(Nindices);
    int* d_indices = thrust::raw_pointer_cast(indices.data());

    int* pos; cudaMallocManaged(&pos, sizeof(int)); *pos = 0;

    AppendArray::Array<int> index(d_indices, Nindices-1, pos);

    int gridsize, blocksize;
    cudaOccupancyMaxPotentialBlockSize(&gridsize, &blocksize, kernelfind, 0, 0);

    kernelfind<<<gridsize, blocksize>>>(d_input, Ninputs, index);
    cudaDeviceSynchronize();

    for(int i = 0; i < *pos; ++i) {
        int idx = indices[i];
        std::cout << i << " " << idx << "  " << inputs[idx] << std::endl;   
    }
    return 0;
}

$ nvcc -std=c++11 -arch=sm_52 -o append append.cu

$ ./append
0 9999  10000
1 19999  20000
2 29999  30000
3 39999  40000
4 49999  50000
5 69999  70000
6 79999  80000
7 59999  60000
8 89999  90000
9 109999  110000
10 99999  100000
11 119999  120000
12 139999  140000
13 129999  130000
14 149999  150000
15 159999  160000
16 169999  170000
17 189999  190000
18 179999  180000
19 199999  200000
20 209999  210000
21 219999  220000
22 239999  240000
23 249999  250000
24 229999  230000
25 279999  280000
26 269999  270000
27 259999  260000
28 319999  320000
29 329999  330000
30 289999  290000
31 299999  300000
32 339999  340000
33 349999  350000
34 309999  310000
35 359999  360000
36 379999  380000
37 399999  400000
38 409999  410000
39 369999  370000
40 429999  430000
41 419999  420000
42 389999  390000
43 439999  440000
44 459999  460000
45 489999  490000
46 479999  480000
47 449999  450000
48 509999  510000
49 539999  540000
50 469999  470000
51 499999  500000
52 569999  570000
53 549999  550000
54 519999  520000
55 589999  590000
56 529999  530000
57 559999  560000
58 619999  620000
59 579999  580000
60 629999  630000
61 669999  670000
62 599999  600000
63 609999  610000
64 699999  700000
65 639999  640000
66 649999  650000
67 719999  720000
68 659999  660000
69 679999  680000
70 749999  750000
71 709999  710000
72 689999  690000
73 729999  730000
74 779999  780000
75 799999  800000
76 809999  810000
77 739999  740000
78 849999  850000
79 759999  760000
80 829999  830000
81 789999  790000
82 769999  770000
83 859999  860000
84 889999  890000
85 879999  880000
86 819999  820000
87 929999  930000
88 869999  870000
89 839999  840000
90 909999  910000
91 939999  940000
92 969999  970000
93 899999  900000
94 979999  980000
95 959999  960000
96 949999  950000
97 1019999  1020000
98 1009999  1010000
99 989999  990000
100 1029999  1030000
101 919999  920000
102 1039999  1040000
103 999999  1000000

推荐阅读