我写了一个简短的CUDA程序,该程序使用高度优化一个旧的四核Intel Q6600处理器(据说所有四个都有〜30 gflops/sec)可以在100,000个元素上进行包含的扫描(或累积/前缀总和),而不是NVIDIA 750 TI(据称能够使用) 1306 gflops/sec的单精度).为什么是这样? 源代码是: #include "cuda_runtime.h" #include "device_launch_parameters.h" #include #include #include #include #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } inline void gpuAssert(cudaError_t code, const char *file, int line, boo
以下是关于 cub 的编程技术问答
我正在CUDA实现中位过滤器.对于特定的像素,我提取与像素周围窗口相对应的邻居,例如N x N(3 x 3)窗口,现在具有N x N元素的数组.我不会为我的应用程序使用超过10 x 10元素的窗口. 此数组现在本地存在于内核中,并且已经加载到设备内存中.从我阅读的以前的SO帖子中,最常见的排序算法是通过推力实现的.但是,只能从主机调用推力.线程 - 用户书面内核 是否有一种快速而有效的方法来对内核内部的一小部分N x N元素进行分类? 解决方案 如果元素的数量固定且较小,则可以使用分类网络(> http://pages.ripco.net/~jgamble/nw.html ).它为固定数量的元素提供了固定数量的比较/交换操作(例如,8个元素的比较/交换迭代). 其他解决方案 您的问题是对许多小阵列进行排序 在罗伯特在评论中提出的建议之后,这个想法是将要排序的小阵列分配给不同的线程块,然后使用让我最终注意到,您的陈述说,cuda推力是在内核中不可召唤的,不
我正在尝试编写一个函数,该函数占据了一个未分类的键/值对,例如 并通过键对其进行分类,同时用相同的键减少对的值: 当前,我使用的是下面的__device__函数,它本质上是一种比特式的排序,它将结合相同键的值并将旧数据设置为无限大的值),以便随后的bitonic排序将它们筛选到底部,并按照int *的值切割的数组. __device__ void interBitonicSortReduce(int2 *sdata, int tid, int recordNum, int *removed) { int n = MIN(DEFAULT_DIMBLOCK, recordNum); for (int k = 2; k
使用CUB :: BlockradixSort在块中进行排序时,如果元素数太大,我们该如何处理?如果我们将瓷砖大小设置为太大,则临时存储的共享内存很快将无法持有.如果我们将其分成多个图块,我们如何在对每个瓷砖进行分类后如何进行后处理? 解决方案 警告:我不是幼崽专家(远非如此). 您可能想查看此 ateys atsya在我在那里做的一些工作的基础. 当然,如果问题大小足够大,则似乎是您可能要考虑的东西.但是您的问题似乎集中于块分类. 从我的测试中,幼崽实际上并没有关于原始数据的位置或放置温度存储的位置的要求.因此,一种可能的解决方案就是将您的温度存储放在全局内存中.为了分析这一点,我创建了一个具有3种不同测试用例的代码: 在全局内存中的温度存储中测试一个Cub块排序的版本. 测试由示例改编的Cub Block排序的原始版本测试从我以前的答案中得出的CUB块排序的版本,该版本没有副本访问/从全局内存中复制.假定数据已经是居民"芯片",即在共享内存中. 所有这
我正在阅读幼崽文档和示例: #include // or equivalently __global__ void ExampleKernel(...) { // Specialize BlockRadixSort for 128 threads owning 4 integer items each typedef cub::BlockRadixSort BlockRadixSort; // Allocate shared memory for BlockRadixSort __shared__ typename BlockRadixSort::TempStorage temp_storage; // Obtain a segment of consecutive items that are blocked across threa
我正在尝试减少库达,而我确实是新手.我目前正在研究Nvidia的示例代码. 我想我真的不确定如何设置块大小和网格大小,尤其是当我的输入数组大(512 X 512)大于单个块大小时. 这是代码. template __global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) { extern __shared__ int sdata[]; unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*(blockSize*2) + tid; unsigned int gridSize = blockSize*2*gridDim.x; sdata[tid] = 0; while (i
我试图在我的“旧"推力代码中引入一些 CUB,因此从一个小示例开始比较 thrust::reduce_by_key 和 cub::DeviceReduce::ReduceByKey,两者都适用于 thrust::device_vectors. 代码的thrust 部分很好,但是CUB 部分天真地使用通过thrust::raw_pointer_cast 获得的原始指针,在CUB 调用后崩溃.我放了一个 cudaDeviceSynchronize() 来尝试解决这个问题,但它没有帮助.CUB 部分代码抄自 CUB 网页. 在 OSX 上,运行时错误是: libc++abi.dylib: terminate called throwing an exception Abort trap: 6 在 Linux 上,运行时错误是: terminate called after throwing an instance of 'thrust::system::system_er