似乎达到了 CUDA 限制，但那是啥限制？

Posted 2023-03-23

【中文标题】似乎达到了 CUDA 限制，但那是啥限制？

【英文标题】：CUDA limit seems to be reached, but what limit is that?似乎达到了 CUDA 限制，但那是什么限制？

【发布时间】：2011-10-18 07:09:18

【问题描述】：

我有一个 CUDA 程序，它似乎达到了某种资源的某种限制，但我不知道该资源是什么。这是核函数：

__global__ void DoCheck(float2* points, int* segmentToPolylineIndexMap, 
                        int segmentCount, int* output)

    int segmentIndex = threadIdx.x + blockIdx.x * blockDim.x;
    int pointCount = segmentCount + 1;

    if(segmentIndex >= segmentCount)
        return;

    int polylineIndex = segmentToPolylineIndexMap[segmentIndex];
    int result = 0;
    if(polylineIndex >= 0)
    
        float2 p1 = points[segmentIndex];
        float2 p2 = points[segmentIndex+1];
        float2 A = p2;
        float2 a;
        a.x = p2.x - p1.x;
        a.y = p2.y - p1.y;

        for(int i = segmentIndex+2; i < segmentCount; i++)
        
            int currentPolylineIndex = segmentToPolylineIndexMap[i];

            // if not a different segment within out polyline and
            // not a fake segment
            bool isLegit = (currentPolylineIndex != polylineIndex && 
                currentPolylineIndex >= 0);      

            float2 p3 = points[i];
            float2 p4 = points[i+1];
            float2 B = p4;
            float2 b;
            b.x = p4.x - p3.x;
            b.y = p4.y - p3.y;

            float2 c;
            c.x = B.x - A.x;
            c.y = B.y - A.y;

            float2 b_perp;
            b_perp.x = -b.y;
            b_perp.y = b.x;

            float numerator = dot(b_perp, c);
            float denominator = dot(b_perp, a);
            bool isParallel = (denominator == 0.0);

            float quotient = numerator / denominator;
            float2 intersectionPoint;
            intersectionPoint.x = quotient * a.x + A.x;
            intersectionPoint.y = quotient * a.y + A.y;

            result = result | (isLegit && !isParallel && 
                intersectionPoint.x > min(p1.x, p2.x) && 
                intersectionPoint.x > min(p3.x, p4.x) && 
                intersectionPoint.x < max(p1.x, p2.x) && 
                intersectionPoint.x < max(p3.x, p4.x) && 
                intersectionPoint.y > min(p1.y, p2.y) && 
                intersectionPoint.y > min(p3.y, p4.y) && 
                intersectionPoint.y < max(p1.y, p2.y) && 
                intersectionPoint.y < max(p3.y, p4.y));
        
    

    output[segmentIndex] = result;

这里是执行内核函数的调用：

DoCheck<<<702, 32>>>(
    (float2*)devicePoints, 
    deviceSegmentsToPolylineIndexMap, 
    numSegments, 
    deviceOutput);

参数大小如下：

devicePoints = 22,464 float2s = 179,712 字节 deviceSegmentsToPolylineIndexMap = 22,463 个整数 = 89,852 个字节 numSegments = 1 int = 4 字节 deviceOutput = 22,463 个整数 = 89,852 个字节

当我执行这个内核时，它会导致显卡崩溃。看来我遇到了某种限制，因为如果我使用DoCheck<<<300, 32>>>(...); 执行内核，它就可以工作。明确一点，参数是一样的，只是块数不一样。

知道为什么一个会导致视频驱动程序崩溃，而另一个不会吗？失败的似乎仍在卡片的块数限制内。

更新 有关我的系统配置的更多信息：

显卡：nVidia 8800GT CUDA 版本：1.1 操作系统：Windows Server 2008 R2

我也在一台笔记本电脑上试了一下，配置如下，结果一样：

显卡：nVidia Quadro FX 880M CUDA 版本：1.2 操作系统：Windows 7 64 位

【问题讨论】：

如果这是显示卡，可能是挂钟时间。显示驱动程序有一个看门狗定时器，它会杀死需要几秒钟才能完成的内核。实施细节和变通方法是特定于操作系统的。您使用的是什么操作系统、卡和 CUDA 版本？

有趣。好的，我将使用该信息更新问题。

看门狗定时器在 Windows 上仍然是一个问题吗？如果是这样，您的内核可能执行时间过长。

那些 CUDA 版本是您卡的计算能力，而不是您使用的 CUDA 版本...但是您肯定会达到显示驱动程序看门狗计时器的限制 - 我认为您正在获得“执行时驱动程序崩溃并被重置”消息？

@Eric：是的，除非使用带有非 WDDM 计算驱动程序的特斯拉卡。

【参考方案1】：

耗尽的资源是时间。在所有当前的 CUDA 平台上，显示驱动程序都包含一个看门狗计时器，它将杀死任何需要超过几秒钟才能执行的内核。在运行显示的卡上运行代码受此限制。

在您使用的 WDDM Windows 平台上，有三种可能的解决方案/解决方法：

弄个Telsa卡，使用TCC驱动，彻底解决问题 尝试修改注册表设置以增加计时器限制（谷歌 TdrDelay 注册表项以获取更多信息，但我不是 Windows 用户，不能比这更具体） 将内核代码修改为“可重入”，并在多个内核启动而不是一次启动中处理数据并行工作负载。内核启动开销并没有那么大，而且处理多个内核运行的工作负载通常很容易实现，具体取决于您使用的算法。

【讨论】：

对于解决方案 #2，这似乎是一个不错的页面：msdn.microsoft.com/en-us/windows/hardware/gg487368

还有一个问题……那么超时对于 CUDA 有什么影响？整个内核调用？每块？每个线程？

任何在看门狗限制内不将 GPU 交给显示驱动程序的主机 API 函数都将触发驱动程序重置事件。在您的情况下，这意味着整个内核（从技术上讲，在最近的 CUDA 版本中的 WDDM 平台上，它也可能意味着任何其他与内核一起批处理的 API 操作）。

Quadro 卡也可以在 TCC 中工作，如果它们没有连接到显示器。