CUDA 和寻址位并行答案

【问题标题】：CUDA and addressing bits in parallelCUDA 和寻址位并行
【发布时间】：2016-07-10 22:06:56
【问题描述】：

我想编写一个 CUDA 程序，它返回包含特定条件的更大数组的位置。

最简单的方法是编写一个内核，它返回一个整数数组，如果满足条件，则返回 1，否则返回 0。

另一种方法可能是只返回找到的索引 - 但根据我对 GPU 同步的了解，这将是有问题的（相当于在 GPU 上实现队列/链表）。

提出的第一个想法的问题是数组将在输入大小内。

我想到的另一种方法是创建一个大小为 log(n)/8+1（n=我检查的项目数）的数组，并为每个数组位置使用 1 位（保存一种压缩表示输出）。

我唯一找不到的是 CUDA 是否支持并行位寻址..

我现在如何做的一个例子：

__global__ void test_kernel(char *gpu, char *gpuFind, int *gputSize, int *gputSearchSize, int *resultsGPU)
{
   int start_idx = threadIdx.x + (blockIdx.x * blockDim.x);

   if (start_idx > *gputTextSize - *gputSearchSize){return;}

   unsigned int wrong=0;
   for(int i=0; i<*gputSearchSize;i++){
     wrong = calculationOnGpu(gpuText, gpuFind, start_idx,i, gputSearchSize);
   }

   resultsGPU[start_idx] = !wrong;

}

我想要做的不是使用 int 或 char 作为“resultsGpu”变量，而是使用其他东西。

谢谢

【问题讨论】：

标签： cuda gpgpu gpu

【解决方案1】：

CUDA GPU can access items 在 1、2、4、8 或 16 字节的边界上。它不具备独立访问字节中位的能力。

一个字节中的位可以通过读取更大的项目来修改，例如char 或int，修改寄存器中的位，然后将该项目写回内存。因此这将是一个读-修改-写操作。

为了在这种多线程场景中保留相邻位，有必要以原子方式更新项目（char、int 等）。没有对 char 数量进行操作的原子，因此需要将这些位分组为 32 的数量，并写入例如作为int。按照这个习惯，每个线程都会做一个原子操作。

32 目前也恰好是经线大小，因此基于经线的内在函数可能是更有效的方法，特别是 warp vote __ballot() 函数。像这样的：

__global__ void test_kernel(char *gpu, char *gpuFind, int *gputSize, int *gputSearchSize, int *resultsGPU)
{
   int start_idx = threadIdx.x + (blockIdx.x * blockDim.x);

   if (start_idx > *gputTextSize - *gputSearchSize){return;}

   unsigned int wrong=0;
     wrong = calculationOnGpu(gpuText, gpuFind, start_idx,0, gputSearchSize);
     wrong = __ballot(wrong);
   if ((threadIdx.x & 31) == 0)
     resultsGPU[start_idx/32] = wrong;

}

您还没有提供完整的代码，所以上面只是一个如何完成的草图。无论如何，我不确定原始内核中的循环是否是一种有效的方法，并且上面假设每个要搜索的数据项有 1 个线程。 __ballot() 应该是安全的，即使在被搜索的数组的一端或另一端存在非活动线程时也是如此。

【讨论】：