我正在编写一个程序,它尝试32位数的每一个组合,查看它是否满足某些条件,并返回那些符合条件的条件。从示例程序中,我一直看到数组的大小总是(No of elements * size of ())。
这个数字看起来太大了,而且大部分的数字也会被拒绝,所以我不需要一个2^32数组。我知道结果的数量将大大少于2^32,但我不知道确切的数量会有多少。
另外,每个线程在尝试数字时都是循环的,所以一个线程有可能有一个以上的正结果。
那么,如何进行内存分配,以及如何存储接受的值?
发布于 2016-03-13 06:55:04
一种方法是尝试分配尽可能多的内存,或者认为需要存储内核输出,然后使用原子递增计数器来跟踪输出缓冲区中的下一个空闲位置,任何给定线程都可以在其中存储结果。
例如,如果您定义了类似于以下内容的助手结构:
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);当内核的输出相当“稀疏”时,这个解决方案可能是最有用的,即输出的数量相对于线程的数量或输入的数量相当小。如果您的问题更“密集”,即内核会相对于线程或输入的数量发出大量的输出,那么原子内存事务可能会造成很大的性能损失。在这种情况下,最好让线程存储到“稀疏”输出缓冲区中,然后使用流压缩传递来消除内核输出缓冲区中的少量空条目。
https://stackoverflow.com/questions/35967086
复制相似问题