在Cuda中实现最大Reduce
我一直在学习Cuda，我仍然在掌握并行性。我现在遇到的问题是在一个值数组上实现一个最大的reduce。这是我的内核
__global__ void max_reduce(const float* const d_array,
float* d_max,
const size_t elements)
{
extern __shared__ float shared[];
int tid = threadIdx.x;
int gid = (blockDim.x * blockIdx.x) + tid;
if (gid < elements)
shared[tid] = d_array[gid];
__syncthreads();
for (unsigned int s=blockDim.x/2; s>0; s>>=1)
{
if (tid < s && gid < elements)
shared[tid] = max(shared[tid], shared[tid + s]);
__syncthreads();
}
if (gid == 0)
*d_max = shared[tid];
}我已经使用相同的方法实现了一个min reduce (用min替换max函数)，它工作得很好。
为了测试内核，我使用一个串行for循环找到了最小和最大值。在内核中，最小值和最大值总是相同的，但只有最小减少值匹配。
有没有明显的我遗漏了什么/做错了什么？
回答 2
Stack Overflow用户
发布于 2013-06-29 23:40:36
您在删除的答案中的主要结论是正确的:您发布的内核没有理解这样一个事实，即在内核执行结束时，您已经完成了大量的总体缩减，但结果并不完全。每个块的结果必须组合在一起(以某种方式)。正如评论中指出的，您的代码还存在一些其他问题。让我们来看看它的一个修改版本：
__device__ float atomicMaxf(float* address, float val)
{
int *address_as_int =(int*)address;
int old = *address_as_int, assumed;
while (val > __int_as_float(old)) {
assumed = old;
old = atomicCAS(address_as_int, assumed,
__float_as_int(val));
}
return __int_as_float(old);
}
__global__ void max_reduce(const float* const d_array, float* d_max,
const size_t elements)
{
extern __shared__ float shared[];
int tid = threadIdx.x;
int gid = (blockDim.x * blockIdx.x) + tid;
shared[tid] = -FLOAT_MAX; // 1
if (gid < elements)
shared[tid] = d_array[gid];
__syncthreads();
for (unsigned int s=blockDim.x/2; s>0; s>>=1)
{
if (tid < s && gid < elements)
shared[tid] = max(shared[tid], shared[tid + s]); // 2
__syncthreads();
}
// what to do now?
// option 1: save block result and launch another kernel
if (tid == 0)
d_max[blockIdx.x] = shared[tid]; // 3
// option 2: use atomics
if (tid == 0)
atomicMaxf(d_max, shared[0]);
}正如Pavan所指出的，您需要初始化您的共享内存阵列。最后启动的块可能不是“完整”块，如果
大于
..。
请注意，在这一行中，即使我们正在检查正在运行的线程(
)小于
当我们添加
至
用于索引到共享内存中
在最后一个块中，我们仍然可以在复制到共享内存中的正当值之外进行索引。因此，我们需要在注释1中指出的共享内存初始化。
正如您已经发现的，您的最后一行是不正确的。每个块都会产生自己的结果，我们必须以某种方式将它们结合起来。如果启动的块数量很少(稍后将对此进行更多介绍)，您可以考虑的一种方法是使用
原子学
..。通常我们会引导人们远离原子，因为它们在执行时间上是“昂贵的”。然而，我们面临的另一个选择是将块结果保存在全局内存中，完成内核，然后可能启动另一个内核来组合各个块结果。如果我最初启动了大量的块(比方说超过1024个)，那么如果我遵循这种方法，我可能最终会启动
两个
附加内核。因此考虑原子学。如上所述，没有本机
函数，但如
文档
，您可以使用
来生成任意的原子函数，我在
，它提供了一个原子最大值，用于
..。
但是运行1024或更多原子函数(每个块一个)是最好的方式吗？可能不会。
当启动线程块的内核时，我们真的只需要启动足够的线程块来使机器保持忙碌。作为一个经验法则，我们希望每个SM至少运行4-8个翘曲，稍微多一点可能是一个好主意。但是，从机器利用率的角度来看，最初启动数千个线程块并没有什么特别的好处。如果我们选择一个数字，比如每个SM有8个线程块，而我们的GPU中最多有14-16个SMs，那么我们得到的是相对较少的8个
*
14 = 112个线程块。让我们选择128 (8
*
16)一个很好的整数。这并没有什么神奇之处，它只是足够让GPU保持忙碌。如果我们让这128个线程块中的每一个都做额外的工作来解决
整体
问题，然后我们可以利用我们对原子的使用，而不会为此付出太多的代价，并避免多次内核启动。那么这看起来会是什么样子呢？
__device__ float atomicMaxf(float* address, float val)
{
int *address_as_int =(int*)address;
int old = *address_as_int, assumed;
while (val > __int_as_float(old)) {
assumed = old;
old = atomicCAS(address_as_int, assumed,
__float_as_int(val));
}
return __int_as_float(old);
}
__global__ void max_reduce(const float* const d_array, float* d_max,
const size_t elements)
{
extern __shared__ float shared[];
int tid = threadIdx.x;
int gid = (blockDim.x * blockIdx.x) + tid;
shared[tid] = -FLOAT_MAX;
while (gid < elements) {
shared[tid] = max(shared[tid], d_array[gid]);
gid += gridDim.x*blockDim.x;
}
__syncthreads();
gid = (blockDim.x * blockIdx.x) + tid; // 1
for (unsigned int s=blockDim.x/2; s>0; s>>=1)
{
if (tid < s && gid < elements)
shared[tid] = max(shared[tid], shared[tid + s]);
__syncthreads();
}
if (tid == 0)
atomicMaxf(d_max, shared[0]);
}使用这个修改过的内核，在创建内核启动时，我们不会根据总体数据大小来决定要启动多少个线程块(
)。相反，我们启动了固定数量的块(例如，128，你可以修改这个数字来找出哪个运行得最快)，并让每个线程块(以及整个网格)在内存中循环，在共享内存中的每个元素上计算部分最大操作。然后，在标记为注释1的行中，我们必须重新设置
变量设置为它的初始值。这实际上是不必要的，如果我们保证网格的大小(
)小于
，这在内核启动时并不难做到。
请注意，在使用此原子方法时，必须初始化结果(
在本例中)设置为适当的值，如
..。
同样，我们通常引导人们远离原子使用，但在这种情况下，如果我们仔细管理它，它是值得考虑的，它允许我们节省额外的内核启动的开销。
有关如何进行快速并行缩减的忍者级分析，请查看Mark Harris的优秀白皮书，该白皮书提供了相关的
CUDA样本
..。
Stack Overflow用户
发布于 2021-02-28 21:43:57
这里有一个看起来很幼稚但并不幼稚的函数。
，但它非常适用于
和
..。
__device__ const float float_min = -3.402e+38;
__global__ void maxKernel(float* d_data)
{
// compute max over all threads, store max in d_data[0]
int i = threadIdx.x;
__shared__ float max_value;
if (i == 0) max_value = float_min;
float v = d_data[i];
__syncthreads();
while (max_value < v) max_value = v;
__syncthreads();
if (i == 0) d_data[0] = max_value;
}是的，没错，只在初始化后同步一次，并在写入结果之前同步一次。该死的比赛条件！全速前进！
在你告诉我它行不通之前，请先试一试。我已经进行了彻底的测试，它每次都能在各种任意大小的内核上工作。事实证明，在这种情况下，竞争条件并不重要，因为while循环解决了这个问题。
它的工作速度比传统的缩减要快得多。另一个令人惊讶的是，内核大小为32的平均通过次数是4。是的，这是(log(n)-1)，这似乎是违反直觉的。这是因为比赛条件给了一个好运气的机会。除了消除传统减少的开销之外，这个额外的好处也随之而来。
对于较大的n，无法避免每个warp至少一次迭代，但该迭代仅涉及一次比较操作，当max为max时，该比较操作通常在warp上立即为false
_
价值处于分布的高端。您可以修改它以使用多个SM，但这将极大地增加总工作负载并增加通信成本，因此不太可能有帮助。
为简洁起见，我省略了大小和输出参数。Size仅仅是线程的数量(可以是137或您喜欢的任何值)。输出返回为
..。
我已经在这里上传了工作文件：
https://github.com/kenseehart/YAMR
https://stackoverflow.com/questions/17371275
复制相似问题
腾讯云开发者
