CUDA 和寻址位并行

Question

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

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

另一种方法可能是 return 仅找到找到的索引 - 但根据我对 GPU 同步的了解，这会产生问题（相当于在 GPU 上实现 queue/linked 列表）。

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

我想到的另一种方法是创建一个大小为 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" 变量，而是使用其他东西。

谢谢

Answer 1

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

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

为了在多线程的情况下保留相邻位，有必要以原子方式更新项目（char、int 等）。没有原子操作在 char 个数量上，因此需要将这些位分组为 32 个数量，并写成例如作为 int。按照这个习惯用法，每个线程都将执行原子操作。

32 也恰好是当前的 warp 大小，因此基于 warp 的内在函数可能是一种更有效的方法，特别是 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;

}

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

CUDA 和寻址位并行

CUDA and addressing bits in parallel

cuda

gpu

gpgpu