0

问题描述

我试图让一个内核总结一个数组的所有元素来工作。内核旨在以每个块 256 个线程和任意数量的块启动。传入的数组长度a总是512的倍数，实际上是#blocks * 512。内核的一个块应该总结“它的”512个元素（256个线程可以使用这个算法总结512个元素），将结果存储在out[blockIdx.x]. 中的值的最终总和out以及块的结果将在主机上完成。
该内核最多可用于 6 个块，即最多 3072 个元素。但是以超过 6 个块启动它会导致第一个块计算出比其他块更大的错误结果（即out = {572, 512, 512, 512, 512, 512, 512})，这个错误的结果是可重现的，多次执行的错误值是相同的。
我想这意味着我的代码中某处存在结构错误，这与blockIdx.x，但唯一的用途是计算blockStart，并且这似乎是正确的计算，对于第一个块也是如此。
我验证了我的主机代码是否为内核计算了正确的块数并传入了一个正确大小的数组。那不是问题。
当然，我在 stackoverflow 上阅读了很多类似的问题，但似乎没有一个描述我的问题（参见此处或此处）
内核是通过managedCuda（C#）调用的，我不知道这是否是一个问题。

硬件

我使用具有以下规格的 MX150：

  • 修订号：6.1
  • 全局内存总量：2147483648
  • 每个块的总共享内存：49152
  • 每个块的总寄存器数：65536
  • 经纱尺寸：32
  • 每个块的最大线程数：1024
  • 最大块数：2147483648
  • 多处理器数量：3

代码

核心

__global__ void Vector_Reduce_As_Sum_Kernel(float* out, float* a)
{   
int tid = threadIdx.x;
int blockStart = blockDim.x * blockIdx.x * 2;
int i = tid + blockStart;

int leftSumElementIdx =  blockStart + tid * 2;

a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];

__syncthreads();

if (tid < 128) 
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if(tid < 64)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid < 32)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid < 16)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid < 8)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid < 4)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid < 2)
{
    a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];
}

__syncthreads();

if (tid == 0)
{
    out[blockIdx.x] = a[blockStart] + a[blockStart + 1];
}
}

内核调用

//Get the cuda kernel
//PathToPtx and MangledKernelName must be replaced
CudaContext cntxt = new CudaContext();
CUmodule module = cntxt.LoadModule("pathToPtx");    
CudaKernel vectorReduceAsSumKernel = new CudaKernel("MangledKernelName", module, cntxt);

//Get an array to reduce
float[] array = new float[4096];
for(int i = 0; i < array.Length; i++)
{
    array[i] = 1;
}

//Calculate execution info for the kernel
int threadsPerBlock = 256;
int numOfBlocks = array.Length / (threadsPerBlock * 2);

//Memory on the device
CudaDeviceVariable<float> m_d = array;
CudaDeviceVariable<float> out_d = new CudaDeviceVariable<float>(numOfBlocks);

//Give the kernel necessary execution info
vectorReduceAsSumKernel.BlockDimensions = threadsPerBlock;
vectorReduceAsSumKernel.GridDimensions = numOfBlocks;

//Run the kernel on the device
vectorReduceAsSumKernel.Run(out_d.DevicePointer, m_d.DevicePointer);

//Fetch the result
float[] out_h = out_d;

//Sum up the partial sums on the cpu
float sum = 0;
for(int i = 0; i < out_h.Length; i++)
{
    sum += out_h[i];
}

//Verify the correctness
if(sum != 4096)
{
    throw new Exception("Thats the wrong result!");
}

更新：

非常有帮助且唯一的答案确实解决了我所有的问题。谢谢！问题是无法预料的比赛条件。

重要提示：

managedCuda 的作者在评论中指出，所有 NPP 方法确实已经在 managedCuda ( using ManagedCuda.NPP.NPPsExtensions;) 中实现。我没有意识到这一点，我想很多人都在阅读这个问题。

4

1 回答 1

1

您没有正确地将每个块将处理整个数组中的 512 个元素的想法合并到您的代码中。根据我的测试，您至少需要进行 2 处更改才能解决此问题：

  1. 在内核中，您错误地计算了每个块的起点：

    int blockStart = blockDim.x * blockIdx.x;

    因为blockDim.x是 256，但每个块处理 512 个元素，您必须将其乘以 2。（在您的计算中乘以 2leftSumElementIdx并没有考虑到这一点——因为它只是乘以tid）。

  2. 在您的主机代码中，您的块数计算不正确：

    vectorReduceAsSumKernel.GridDimensions = array.Length / threadsPerBlock;

    对于 2048array.Length的值和 256 的值threadsPerBlock，这将创建 8 个块。但正如您已经指出的那样，您的意图是启动区块（2048/512）。所以你需要将分母乘以 2：

    vectorReduceAsSumKernel.GridDimensions = array.Length / (2*threadsPerBlock);

此外，您的减少扫描模式被破坏了。它依赖于warp-execution-order，以给出正确的结果，并且CUDA没有指定warp执行顺序。

要了解原因，让我们举一个简单的例子。让我们只考虑一个线程块，数组的起点全为 1，就像您初始化它一样。

现在，warp 0 由线程 0-31 组成。你的缩减扫描操作是这样的：

a[i] = a[leftSumElementIdx] + a[leftSumElementIdx + 1];

所以warp 0中的每个线程都会收集另外两个值并将它们相加，然后存储它们。线程 31 将获取这些值a[62]并将a[63]它们相加。如果 和 的值a[62]仍然a[63]是 1，就像初始化的那样，那么这将按预期工作。a[62]但是和的值由 warp 1a[63] 写入，由线程 32-63 组成。因此，如果 warp 1 在 warp 0 之前执行（完全合法），那么您将得到不同的结果。这是一个全局内存竞争条件。这是因为您的输入数组既是中间结果的来源又是目的地，并且__syncthreads()不会为您解决这个问题。它不会强制扭曲以任何特定的顺序执行。

一种可能的解决方案是修复您的扫描模式。在任何给定的归约循环中，让我们有一个扫描模式，其中每个线程写入和读取在该循环期间没有被任何其他线程触及的值。您的内核代码的以下改编实现了这一点：

__global__ void Vector_Reduce_As_Sum_Kernel(float* out, float* a)
{
  int tid = threadIdx.x;
  int blockStart = blockDim.x * blockIdx.x * 2;
  int i = tid + blockStart;

  for (int j = blockDim.x; j > 0; j>>=1){
    if (tid < j)
      a[i] += a[i+j];

    __syncthreads();}

  if (tid == 0)
  {
    out[blockIdx.x] = a[i];
  }
}

对于通用减少，这仍然是一种非常缓慢的方法。本教程介绍了如何编写更快的归约。而且，正如已经指出的那样，managedCuda 可能有一些方法可以完全避免编写内核。

于 2018-04-15T21:30:07.790 回答