我正在编写一个程序，它尝试 32 位数字的每个组合，查看它是否满足某些条件并返回满足的条件。从示例程序中，我看到数组的大小始终为 (No of elements * sizeof())。
这个数字似乎太大了，而且大多数数字也会被拒绝，所以我不需要 2^32 数组。我知道结果的数量将大大少于 2^32，但我也不知道究竟会有多少。
此外，每个线程在尝试数字时都会循环，因此一个线程可能有多个阳性结果。
那么如何进行内存分配以及如何存储接受的值？
【问题讨论】：
一种方法是尝试分配尽可能多的内存，或者认为需要存储内核输出，然后使用原子递增的计数器来跟踪输出缓冲区中任何给定线程可以存储的下一个空闲位置结果。
例如，如果你定义一个类似这样的辅助结构：
struct counter
{
unsigned int * _val;
__host__ __device__
counter(unsigned int * value) : _val(value) {};
__device__
unsigned int next() {
return atomicAdd(_val, 1);
};
}
然后在主机代码中执行类似的操作
unsigned int * array_index;
const unsigned int zero = 0;
cudaMalloc((void **)&array_index, sizeof(unsigned int*));
cudaMemcpy(array_index, &zero, sizeof(unsigned int), cudaMemcpyHostToDevice);
counter mycounter(array_index);
您有一个零初始化的设备内存计数器，可以通过重复调用next() 方法在设备代码中安全地读取和递增该计数器。
在内核中是这样的：
__global__ void kernel(Type * buffer, counter mycounter)
{
// Calculate and find a match...
buffer[mycounter.next()] = match;
}
[强烈警告：所有在浏览器中编写的代码，未经编译或测试，可能会使您的 GPU 着火，使用风险自负]
然后，您的内核可以为每个线程发出尽可能多的输出，以适合您的算法设计。扩展我上面说明的设计模式以包括对数组的边界检查是明智的。您还应该注意内核发出的输出总数可以这样检索：
unsigned int N;
cudaMemcpy(&N, array_index, sizeof(unsigned int), cudaMemcpyDeviceToHost);
当内核的输出相当“稀疏”时，此解决方案可能最有用，即输出数量相对于线程数量或输入数量相当小。如果您的问题更“密集”，即内核将发出大量相对于线程数或输入数的输出，那么原子内存事务可能会导致显着的性能损失。在这种情况下，最好将线程存储到“稀疏”输出缓冲区中，然后使用流压缩传递来消除内核输出缓冲区中的少量空条目。
【讨论】：