我将以下代码作为重组数据的一部分,以便以后在CUDA内核中使用:
thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(dev_particle_cell_indices);
int total = 0;
for(int i = 0; i < num_cells; i++) {
particle_offsets[i] = total;
// int num = 0;
int num = thrust::count(dev_ptr, dev_ptr + num_particles, i);
particle_counts[i] = num;
total += num;
}
现在,如果我设置num为0(取消注释第5行,并注释掉第6行),应用程序将以超过30 fps的速度运行,这是我的目标.但是,当我设置num等于thrust::count呼叫时,帧率降至约1-2 fps.为什么会这样？
我的理解是,推力应该是高度优化的算法的集合,利用GPU的强大功能,所以我很惊讶它会对我的程序的性能产生这样的影响.这是我第一次使用推力,所以我可能没有意识到一些重要的细节.
是否有关于thrust::count在循环中使用导致它运行如此缓慢的事情？如何优化我的使用？
在我目前的测试案例中,给出一些数字num_particles大约是2000,num_cells大概是1500.
我必须坦率地说,这将是
对推力的批评
他们在优化矢量输入的并行算法方面做得很好.他们使用的数据级并行(除其他事项外),以parllelize的工作真的很好的算法大,矢量投入.但是他们没有改进它并且一直走向真正的数据级并行性.即大量的小问题.
第二种情况在许多实际应用中都很有用,而ArrayFire在这方面提供了解决方案(看看gfor,并行for循环).
原本应该简单地调用减少和扫描的是4种算法(其中一种算法代价很高)和3种内存副本.
以下是代码在ArrayFire中的工作方式:
array cell_indices(num_particles, 1, dev_particle_cell_indices, afDevicePointer);
array particle_counts = zeros(num_cells);
gfor(array i, num_cells) // Parallel for loop
particle_counts(i) = sum(cell_indices == i);
array particle_offsets = accum(particle_counts); // Inclusive sum
-
建立
基准1
使用num_particles = 2000和num_cells = 1500(就像原始问题一样)
$ ./a.out
Thrust time taken: 0.002384
ArrayFire time taken: 0.000131
ArrayFire是18倍的速度
基准2
使用num_particles = 10000和num_cells = 2000(像talonmies的测试用例)
$ ./a.out
Thrust time taken: 0.002920
ArrayFire time taken: 0.000132
ArrayFire是22倍的速度
基准3
使用num_particles = 50000和num_cells = 5000(只是一个更大的测试用例)
$ ./a.out
Thrust time taken: 0.003596
ArrayFire time taken: 0.000157
ArrayFire是23倍的速度
笔记
结论
如果您可以重新编写问题,那么推力确实可以提供适当的加速.但这并不总是可行的,对于更复杂的问题来说并非易事.这些数字表明存在更高性能的余地(因为数据并行度很高),而这种情况根本没有得到推动.
ArrayFire以更有效的方式利用并行资源,时间表明gpu仍未饱和.
您可能想要编写自己的自定义cuda代码或使用ArrayFire.我只想指出,有时使用推力不是一种选择,因为它在大量的小问题上几乎没用.
编辑修复了基准1的结果(我使用了错误的数字)
性能thrust::count很好,这是你尝试使用它的方式,这对性能有问题.如果你有很多粒子而且只有几个单元格,那么你的实现thrust::count可能并不是一个坏主意.你的问题是你有1500个细胞.这意味着count每次要进行计算时,1500个调用和1500个设备到主机内存传输.正如您所发现的,所有内核启动和所有PCI-e总线副本的延迟都会影响性能.
对于大量细胞更好的方法是这样的:
thrust::device_ptr<int> rawin = thrust::device_pointer_cast(dev_particle_cell_indices);
// Sort a scratch copy of the cell indices by value
thrust::device_vector<int> cidx(num_particles);
thrust::copy(rawin, rawin+num_particles, cidx.begin());
thrust::sort(cidx.begin(), cidx.end());
// Use binary search to extract all the cell counts/offsets
thrust::counting_iterator<int> cellnumber(0);
thrust::device_vector<int> offsets(num_cells), counts(num_cells);
// Offsets come from lower_bound of the ordered cell numbers
thrust::lower_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, offsets.begin());
// Counts come from the adjacent_difference of the upper_bound of the ordered cell numbers
thrust::upper_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, counts.begin());
thrust::adjacent_difference(counts.begin(), counts.end(), counts.begin());
// Copy back to the host pointer
thrust::copy(counts.begin(), counts.end(), particle_counts);
thrust::copy(offsets.begin(), offsets.end(), particle_offsets);
在这里,我们首先对单元索引的本地副本进行排序,然后使用推力二进制搜索函数执行与代码相同的操作,但是通过GPU内存中的数据传递的次数要少得多,而只有两个设备来托管内存副本以获得所有结果都回到主机.
当我thrust::count使用上面发布的代码对我的实现进行基准测试以获得一个非常重要的案例(在OS X上使用CUDA 4.1的GeForce 320M上有10000个随机粒子和2000个单元格),我发现你的版本需要大约0.95秒才能运行,而排序/搜索版本需要大约0.003秒才能运行.因此,如果您使用更有效的策略和更合适的算法,使用推力可能会有几百倍的加速.