Cuda 有效地将数据插入到未排序的填充数组中

Question

我在 Cuda 中有两个数组；

int *main; // unsorted
int *source; // sorted

我的部分算法要求我定期将新数据从源数组插入到主数组中。如果主数组中的某个位置为零，则假定它为空，因此可以使用源数组中的值填充它。

我只是想知道最有效的方法是什么，我尝试了几种方法，但仍然认为这里还有一些性能提升。

目前我正在使用基数排序的修改版本，将主数组的内容“洗牌”到主数组的最后，将所有零值留在数组的开头，从来源琐碎。排序已修改为迭代单个位，而不是 32 位，只需在输入上进行简单切换即可工作；

input[i] = source[i] > 1 ? 1 : 0

我想知道这是否已经是一种非常有效的方法？我想知道我是否不会通过使用战术部署的 atomicAdd 来获得一些东西，例如；

__global__ void find(int *destination, int *indices, const int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if((destination[idx] == 0)&&(count<elements_to_add))
    {
        indices[count] = idx;
        atomicAdd(&count, 1);
    }
}

__global__ void insert(int *destination, int *indices, int *source, const int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if((source[idx] > 0)&&(indices[idx] > 0))
    {
        destination[indices[idx]] = source[idx];
    }
}

find<<<G,T>>>(...);
insert<<<G,T>>>(...);

目前我没有通过源数组插入那么多项目，但将来可能会改变。

感觉这应该是一个以前已经解决的常见问题，我想知道推力库是否有帮助，但是浏览适当的函数后感觉不太适合我想要的完成（不太符合我已有的代码）

感谢经验丰富的 Cuda 开发人员的想法！

Answer 1

您可以将被归类为流压缩过程的查找算法与被归类为分散过程的插入解耦。但是，您可以合并两者的功能。

假设srcPtr 是一个指针，它的内容驻留在全局内存中，并且在内核启动之前已经设置为零。

__global__ void find_and_insert( int* destination, int const* source, int const N, int* srcPtr ) {    // Assuming N is the length of the destination buffer and also the length of the source buffer is less than N.

int const idx = blockIdx.x * blockDim.x + threadIdx.x;

// Get the assigned element.
int const dstElem = destination[ idx ];
bool const pred = ( dstElem == 0 );

// Intra-warp binary reduction to count the total number of lanes with empty elements.
int const predBallot = __ballot( pred );
int const intraWarpRed = __popc( predBallot );

// Warp-aggregated atomics to reduce the contention over the srcPtr content.
unsigned int laneID; asm( "mov.u32 %0, %laneid;" : "=r"(laneID) ); //const uint laneID = tidWithinCTA & ( WARP_SIZE - 1 );
int posW;
if( laneID == 0 )
    posW = atomicAdd( srcPtr, intraWarpRed );
posW = __shfl( posW, 0 );

// Threads that have found empty elements can fill out their assigned positions from the src. Intra-warp binary prefix sum is used here.
uint laneMask; asm( "mov.u32 %0, %lanemask_lt;" : "=r"(laneMask) ); //const uint laneMask =  0xFFFFFFFF >> ( WARP_SIZE - laneID ) ;
int const positionToRead = posW + __popc( predBallot & laneMask );
if( pred )
    destination[ idx ] = source[ positionToRead ];

}

一些事情：

1. 这个内核只是一个关于如何做的建议。在这里，warp 中的线程在任务上进行协作。您可以在线程块上扩展二进制归约和前缀和。
2. 我在浏览器中编写了这个内核，还没有测试过。所以要小心。
3. 整个设计并不是什么新鲜事。已经实现了类似的方法（例如this paper），并且主要基于the work done by Mark Harris and Michael Garland。