C++：简单的CUDA卷重构代码崩溃

Question

我目前正在开发一个涉及CUDA的更全面的项目。在最近几天里，我一直在遇到错误，我一直在拼命地尝试bug修复。但是，我无法弄清楚，所以现在我构成了一个最小的例子，它显示了相同的行为。我不得不说我对CUDA很陌生。我正在使用Visual Studio 2015和CUDA Toolkit 7.5。

该程序涉及在GPU内存上创建3D卷，然后计算值并将其写入卷。我试图使代码尽可能简单：

首先IST的main.cpp文件：

#include "cuda_test.h" 

int main() { 

    size_t const xDimension = 500; 
    size_t const yDimension = 500; 
    size_t const zDimension = 1000; 

    //allocate volume part memory on gpu 
    cudaPitchedPtr volume = ct::cuda::create3dVolumeOnGPU(xDimension, yDimension, zDimension); 

    //start reconstruction 
    ct::cuda::startReconstruction(volume, 
            xDimension, 
            yDimension, 
            zDimension); 

return 0; 

}

然后cuda_test.h这是实际.CU文件头文件：

#ifndef CT_CUDA 
#define CT_CUDA 

#include <cstdlib> 
#include <stdio.h> 
#include <cmath> 

//CUDA 
#include <cuda_runtime.h> 

namespace ct { 

    namespace cuda { 

     cudaPitchedPtr create3dVolumeOnGPU(size_t xSize, size_t ySize, size_t zSize); 
     void startReconstruction(cudaPitchedPtr volume, 
           size_t xSize, 
           size_t ySize, 
           size_t zSize); 

    } 

} 

#endif 

包含行为

然后是cuda_test.cu文件UAL功能实现：

#include "cuda_test.h" 

namespace ct { 

    namespace cuda { 

     cudaPitchedPtr create3dVolumeOnGPU(size_t xSize, size_t ySize, size_t zSize) { 
      cudaExtent extent = make_cudaExtent(xSize * sizeof(float), ySize, zSize); 
      cudaPitchedPtr ptr; 
      cudaMalloc3D(&ptr, extent); 
      printf("malloc3D: %s\n", cudaGetErrorString(cudaGetLastError())); 
      cudaMemset3D(ptr, 0, extent); 
      printf("memset: %s\n", cudaGetErrorString(cudaGetLastError())); 
      return ptr; 
     } 

     __device__ void addToVolumeElement(cudaPitchedPtr volumePtr, size_t ySize, size_t xCoord, size_t yCoord, size_t zCoord, float value) { 
      char* devicePtr = (char*)(volumePtr.ptr); 
      //z * xSize * ySize + y * xSize + x 
      size_t pitch = volumePtr.pitch; 
      size_t slicePitch = pitch * ySize; 
      char* slice = devicePtr + zCoord*slicePitch; 
      float* row = (float*)(slice + yCoord * pitch); 
      row[xCoord] += value; 
     } 

     __global__ void reconstructionKernel(cudaPitchedPtr volumePtr, size_t xSize, size_t ySize, size_t zSize) { 

      size_t xIndex = blockIdx.x; 
      size_t yIndex = blockIdx.y; 
      size_t zIndex = blockIdx.z; 

      if (xIndex == 0 && yIndex == 0 && zIndex == 0) { 
       printf("kernel start\n"); 
      } 

      //just make sure we're inside the volume bounds 
      if (xIndex < xSize && yIndex < ySize && zIndex < zSize) { 

       //float value = z; 
       float value = sqrt(sqrt(sqrt(5.3))) * sqrt(sqrt(sqrt(1.2))) * sqrt(sqrt(sqrt(10.8))) + 501 * 0.125 * 0.786/5.3; 

       addToVolumeElement(volumePtr, ySize, xIndex, yIndex, zIndex, value); 

      } 

      if (xIndex == 0 && yIndex == 0 && zIndex == 0) { 
       printf("kernel end\n"); 
      } 

     } 

     void startReconstruction(cudaPitchedPtr volumePtr, size_t xSize, size_t ySize, size_t zSize) { 
      dim3 blocks(xSize, ySize, zSize); 
      reconstructionKernel <<< blocks, 1 >>>(volumePtr, 
                xSize, 
                ySize, 
                zSize); 
      printf("Kernel launch: %s\n", cudaGetErrorString(cudaGetLastError())); 
      cudaDeviceSynchronize(); 
      printf("Device synchronise: %s\n", cudaGetErrorString(cudaGetLastError())); 
     } 

    } 

} 

功能create3dVolumeOnGPU分配在GPU存储器3维“音量”，并返回一个指向它的指针。这是一个主机功能。第二个主机功能是startReconstruction。它所做的唯一的事情就是启动实际的内核，使用与卷中的体素一样多的块。内核函数是reconstructionKernel。它只是计算一些常数中的任意值，然后调用addToVolumeElement（设备函数）将结果写入相应的体素（添加它）。

现在，问题是它崩溃了。如果我和调试器（NSight）推出，NSight中断给错误消息：

CUDA grid launch failed: CUcontext: 2358451327088 CUmodule: 2358541519888 Function: _ZN2ct4cuda20reconstructionKernelE14cudaPitchedPtryyy

控制台输出：

malloc3D: no error 
memset: no error 
kernel started 
kernel end 

如果我在释放模式启动整个机器复位。

但是，如果我改变体积的尺寸要小一些它的作品，例如：

size_t const xDimension = 100; 
    size_t const yDimension = 100; 
    size_t const zDimension = 100; 

然而，自由GPU内存的数量不应该是问题（卡有4GB VRAM）。

这将是很好，如果有人可以看看它，也许给我一个小费可能会导致问题。现在

Answer 1

，问题是它崩溃

这将是很好，如果有人可以看看它，也许给我一个提示，是什么引发的问题。

我想这很可能是您碰到a WDDM TDR issue。在Windows上，任何时候在WDDM GPU上运行的内核执行时间都需要大约2秒钟，您可能会遇到WDDM TDR看门狗（假设您没有对看门狗进行任何更改）。

此外，启动内核是这样的：

reconstructionKernel <<< blocks, 1 >>>(...); 

其中线程每块数为1，意味着只有一个在各经纱（以及在每个块中）线程是活动的。但GPU喜欢每个warp有32个活动线程。所以净效应是GPU资源的低效利用;也许，当你运行的内核这样的GPU马力高达97％闲置。

所以，如果你的代码是非常灵活，允许这样的：

reconstructionKernel <<< blocks, 1 >>>(...); 

或等价的：

reconstructionKernel <<< blocks/256, 256 >>>(...); 

（这只是一个代表性的例子，我知道你有一个多维网格，上述可能不适合你的情况正是相关）

然后第二次调用方法几乎肯定会更有效率，导致缩短执行时间同样的工作。

所以我相信，当你每块多线程测试你的代码，你不喜欢的东西上面，它低于该TDR限制的执行时间。

这是一个完全正常的解决方案，但如果你最终增加更多的工作，你的内核（更多的总线程，或每个线程更多的工作），那么你可以再次运行到了极限。在这种情况下，链接的文章解释了可能的解决方法。

顺便说一句，内核启动配置是这样的：

kernel<<<1, ?>>>(...); 

或本：

kernel<<<?, 1>>>(...); 

是从未建议在GPU高性能的代码。