CUDA 超时？ /费米/gtx465答案

【问题标题】：Timeout in CUDA? / fermi / gtx465CUDA 超时？ /费米/gtx465
【发布时间】：2012-04-14 16:56:58
【问题描述】：

我在 MS VS2005 上使用 CUDA SDK 3.1，GPU GTX465 1 GB。我有这样一个核函数：

__global__ void CRT_GPU_2(float *A, float *X, float *Y, float *Z, float *pIntensity, float *firstTime, float *pointsNumber)
{


  int holo_x = blockIdx.x*20 + threadIdx.x;
  int holo_y = blockIdx.y*20 + threadIdx.y;

  float k=2.0f*3.14f/0.000000054f;

  if (firstTime[0]==1.0f)
  {
   pIntensity[holo_x+holo_y*MAX_FINAL_X]=0.0f; 
  }

  for (int i=0; i<pointsNumber[0]; i++)
  {
   pIntensity[holo_x+holo_y*MAX_FINAL_X]=pIntensity[holo_x+holo_y*MAX_FINAL_X]+A[i]*cosf(k*sqrtf(pow(holo_x-X[i],2.0f)+pow(holo_y-Y[i],2.0f)+pow(Z[i],2.0f)));
  }

  __syncthreads(); 


}

这是调用内核函数的函数：

extern "C" void go2(float *pDATA, float *X, float *Y, float *Z, float *pIntensity, float *firstTime, float *pointsNumber)
{
 dim3 blockGridRows(MAX_FINAL_X/20,MAX_FINAL_Y/20);
 dim3 threadBlockRows(20, 20);

 CRT_GPU_2<<<blockGridRows, threadBlockRows>>>(pDATA, X, Y, Z, pIntensity,firstTime, pointsNumber); 
 CUT_CHECK_ERROR("multiplyNumbersGPU() execution failed\n");
 CUDA_SAFE_CALL( cudaThreadSynchronize() );
}

我正在循环中加载此函数的所有参数（例如，在一次循环迭代中每个参数有 4096 个元素）。总的来说，我想在所有循环迭代之后为每个参数制作 32768 个元素的内核。

MAX_FINAL_X 是 1920，MAX_FINAL_Y 是 1080。

当我开始算法时，第一次迭代非常快，经过一两次迭代后，我得到了有关 CUDA 超时错误的信息。我在 GPU gtx260 上使用了这个算法，据我记得它做得更好......

你能帮我吗..根据这个算法中的新费米拱门，我可能犯了一些错误吗？

【问题讨论】：

标签： timeout cuda nvidia

【解决方案1】：

最好打电话 CUT_CHECK_ERROR 之后 cudaThreadSynchronize()。因为内核异步运行，您必须等待内核结局知道错误...也许在第二次迭代中你会收到一个错误从第一次使用内核开始。
请务必你在最有趣的变量中有一些有效数字 pointsNumber[0]（可能会导致长的内部循环）。
你也可以提高内核的速度功能：
- 使用更好的方块。线程配置 20x20 会导致非常慢的内存使用（请参阅编程指南和最佳实践）。尝试使用 16x16 的块。
- 不要使用pow(..., 2.0)函数。使用 SQR 宏更快 (#define SQR(x) (x)*(x))
- 您不使用共享内存，因此不需要__syncthreads()。

PS：您还可以将值参数传递给 CUDA 函数，而不仅仅是指针。速度会一样。

PPS：请提高代码的可读性...现在您必须编辑六个位置来更改块配置...在内核中您可以使用blockDim 变量，您可以在go2 函数中使用常量。你也可以使用bool firstTime——它会比float好得多。

【讨论】：

我的 pointNumber[0] 值有误，而且你的线索将这个算法加速了几乎 2 倍 :D 谢谢 :) 我有第二个问题.. 可能是 gpu 的变化（表格 gtx260 896MB到 gtx465 1GB) 可能会导致内存管理出现一些错误？在 gtx260 上，我可以分配 500*500*500 的浮点数组（这是一种 LUT 数组）浮点数，现在在 gtx 上，当我分配超过 60*60*60 的浮点数数组时，出现错误“未指定的启动错误”将数据从主机复制到设备内存（此 memcopy 未与该 LUT 数组连接）...？
没有代码很难说什么......我在memcpy期间看到了这样的错误。此错误与主机程序中的“分段错误”非常相似，并且发生在内核中。 CUT_CHECK_ERROR 读取可能由早期内核启动引起的错误消息（如果您不使用 cudaThreadSynchronize）。异步操作只有 4 种类型： 1. 内核启动 2. 所有 cudaMemcpy*Async 3. Memcpy 设备设备 4. 内存初始化一些建议： - 检查cudaMemcpy 中的参数。正确的顺序和正确的内存分配（也许你忘了cudaMalloc？） - 使用emu模式来dbg。
还有 1 个猜测。您是否将边界检查添加到内核？ if (holo_x+holo_y*MAX_FINAL_X >= MAX_FINAL_X*MAX_FINAL_Y) return; 实际上，2 次 - 不是一个限制）如果你将使用 constant 内存来存储 A、X、Y 和 Z + 可能会转向快速数学（小心）你会给出更多加速。
糟糕！我的错。检查必须是'if (holo_x >= MAX_FINAL_X || holo_y >= MAX_FINAL_Y) return;'
顺便说一句，我提倡避免使用 cutil（即 CUT_CHECK_ERROR）。它不受 NVIDIA 支持，而且它不是处理错误的好方法，因为它会立即退出。程序员应该以适合应用程序的方式明确地处理错误。在 SDK 示例中使用 cutil 宏来隐藏它以专注于实际算法等，但它们不是实际程序的最佳实践。

【解决方案2】：

您的 GPU 是否连接到显示器？如果是这样，我相信默认情况下内核执行将在 5 秒后中止。您可以使用cudaGetDeviceProperties 检查内核执行是否会超时 - 请参阅reference page

【讨论】：

【解决方案3】：

在内核的周期中，您在同一个数组中写入，然后从中读取 - 对于全局内存使用而言，这是最糟糕的，因为来自不同块的扭曲相互等待。

【讨论】：