在 CUDA 内核中循环数据会导致应用程序中止

Question

问题：

当我增加 CUDA kernel 循环内正在处理的数据量时 - 它会导致应用程序中止！

异常：

ManagedCuda.CudaException: 'ErrorLaunchFailed: An exception occurred on the device while executing a kernel. Common causes include dereferencing an invalid device pointer and accessing out of bounds shared memory.

问题：

如果有人能阐明我在当前实施中遇到的限制以及导致应用程序崩溃的确切原因，我将不胜感激。

或者，我附上了完整的内核代码，如果有人能说出如何在不抛出异常的情况下以这种方式重新建模。这个想法是内核接受 combinations 然后对同一组 data 执行计算（在循环中）。因此，内部的循环计算应该是顺序的。内核本身的执行顺序无关紧要。这是组合数学问题。

欢迎任何建议。

代码（简短版本，足以中止应用程序）：

extern "C"
{
    __device__ __constant__ int arraySize;

    __global__ void myKernel(
        unsigned char* __restrict__  output,
        const int* __restrict__  in1,
        const int* __restrict__  in2,
        const double* __restrict__  in3,
        const unsigned char* __restrict__  in4)
    {
        for (int row = 0; row < arraySize; row++)
        {
            // looping over sequential data.
        }
    }
}

在上面的示例中，如果 arraySize 接近 50_000，则应用程序开始中止。使用相同类型的输入参数，如果我们覆盖或 hardcore arraySize 到 10_000 然后代码成功完成。

代码-内核（完整版）

#iclude <cuda.h> 
#include "cuda_runtime.h"
#include <device_launch_parameters.h> 
#include <texture_fetch_functions.h> 
#include <builtin_types.h> 

#define _SIZE_T_DEFINED

#ifndef __CUDACC__
#define __CUDACC__
#endif

#ifndef __cplusplus
#define __cplusplus
#endif

texture<float2, 2> texref;

extern "C"
{
    __device__ __constant__ int width;
    __device__ __constant__ int limit;
    __device__ __constant__ int arraySize;

    __global__ void myKernel(
        unsigned char* __restrict__  output,
        const int* __restrict__  in1,
        const int* __restrict__  in2,
        const double* __restrict__  in3,
        const unsigned char* __restrict__  in4)
    {
        int index = blockIdx.x * blockDim.x + threadIdx.x;

        if (index >= limit)
            return;

        bool isTrue = false;
        int varA = in1[index];
        int varB = in2[index];

        double calculatable = 0;
        for (int row = 0; row < arraySize; row++)
        {
            if (isTrue)
            {
                int idx = width * row + varA;
                if (!in4[idx])
                    continue;

                calculatable = calculatable + in3[row];
                isTrue = false;
            }
            else
            {
                int idx = width * row + varB;
                if (!in4[idx])
                    continue;

                calculatable = calculatable - in3[row];
                isTrue = true;
            }
        }

        if (calculatable >= 0) {
            output[index] = 1;
        }
    }
}

代码-主机（完整版）

    public static void test()
    {
        int N = 10_245_456; // size of an output

        CudaContext cntxt = new CudaContext();
        CUmodule cumodule = cntxt.LoadModule(@"kernel.ptx");
        CudaKernel myKernel = new CudaKernel("myKernel", cumodule, cntxt);

        myKernel.GridDimensions = (N + 255) / 256;
        myKernel.BlockDimensions = Math.Min(N, 256);

        // output
        byte[] out_host = new byte[N]; // i.e. bool
        var out_dev = new CudaDeviceVariable<byte>(out_host.Length);

        // input
        int[] in1_host = new int[N];
        int[] in2_host = new int[N];
        double[] in3_host = new double[50_000]; // change it to 10k and it's OK
        byte[] in4_host = new byte[10_000_000]; // i.e. bool
        var in1_dev = new CudaDeviceVariable<int>(in1_host.Length);
        var in2_dev = new CudaDeviceVariable<int>(in2_host.Length);
        var in3_dev = new CudaDeviceVariable<double>(in3_host.Length);
        var in4_dev = new CudaDeviceVariable<byte>(in4_host.Length);

        // copy input parameters
        in1_dev.CopyToDevice(in1_host);
        in2_dev.CopyToDevice(in2_host);
        in3_dev.CopyToDevice(in3_host);
        in4_dev.CopyToDevice(in4_host);

        myKernel.SetConstantVariable("width", 2);
        myKernel.SetConstantVariable("limit", N);
        myKernel.SetConstantVariable("arraySize", in3_host.Length);

        // exception is thrown here
        myKernel.Run(out_dev.DevicePointer, in1_dev.DevicePointer, in2_dev.DevicePointer,in3_dev.DevicePointer, in4_dev.DevicePointer);

        out_dev.CopyToHost(out_host);
    }

分析

我最初的假设是我有内存问题，但是，根据 VS 调试器，我在主机环境中遇到的数据略高于 500mb。所以我想无论我向 GPU 复制多少数据 - 它都不应该超过 1Gb 甚至最大值 11Gb。后来我注意到只有当内核内部的循环有许多数据记录要处理时才会发生崩溃。这让我相信我遇到了某种线程超时限制或类似的问题。没有确凿的证据。

系统

我的系统规格是 Ram 的 16Gb 和 GeForce 1080 Ti 11Gb。使用 Cuda 9.1. 和 managedCuda 版本 8.0.22（也尝试使用 master 分支的 9.x 版本）

编辑 1：2018 年 4 月 26 日 刚刚测试了相同的逻辑，但仅在 OpenCL 上进行了测试。代码不仅成功完成，而且执行速度比 CUDA 好 1.5-5 倍，具体取决于输入参数大小：

kernel void Test (global bool* output, global const int* in1, global const int* in2, global const double* in3, global const bool* in4, const int width, const int arraySize)
{
    int index = get_global_id(0);

    bool isTrue = false;
    int varA = in1[index];
    int varB = in2[index];

    double calculatable = 0;

    for (int row = 0; row < arraySize; row++)
    {
        if (isTrue)
        {
            int idx = width * row + varA;

            if (!in4[idx]) {
                continue;
            }

            calculatable = calculatable + in3[row];
            isTrue = false;
        }
        else
        {
            int idx = width * row + varB;

            if (!in4[idx]) {
                continue;   
            }

            calculatable = calculatable - in3[row];
            isTrue = true;
        }
    }

    if (calculatable >= 0)
    {
        output[index] = true;
    }
}

我真的不想在这里开始OpenCL/CUDA war。如果我在最初的 CUDA 实施中有任何需要注意的地方 - 请告诉我。

编辑：26.04.2018。在遵循评论部分的建议后，我能够在抛出异常之前将处理的数据量增加 3 倍。我能够通过切换到 Release 模式而不是 Debug 模式生成的 .ptx 来实现这一点。此改进可能与以下事实有关：在 Debug 设置中，我们还将 Generate GPU Debug information 设置为 Yes 以及其他可能影响性能的不必要设置。我现在将尝试搜索有关如何进行的信息内核的计时可以增加。我仍然没有达到 OpenCL 的结果，但已经接近了。

对于 CUDA 文件生成，我正在使用 VS2017 Community、CUDA 9.1 项目、v140 toolset、为 x64 平台构建、post 构建事件已禁用，配置类型：utility。代码生成设置为：compute_30,sm_30。例如，我不确定为什么不是 sm_70。我没有其他选择。

Answer 1

我已经设法将 CUDA 性能提高到 OpenCL。更重要的是，代码现在可以无异常地完成执行。感谢 Robert Crovella。 谢谢！

在显示这里的结果之前有一些规格：

CPUIntel i7 8700k12核(6+6)
GPU GeForce 1080 Ti 11Gb

这是我的结果 (library/technology)：

CPU 并行 for 循环：607907 毫秒（默认）
GPU（Alea、CUDA）：9905 毫秒 (x61)
GPU（managedCuda、CUDA）：6272 毫秒 (x97)
GPU（Coo、OpenCL）：8277 毫秒 (x73)

解决方案 1：

解决方案是将 WDDM TDR Delay 从默认的 2 秒增加到 10 秒。 As easy as that.

解法二：

我能够通过以下方式获得更多性能：

正在将 CUDA 项目属性中的 compute_30,sm_30 设置更新为 compute_61,sm_61
使用 Release 设置而不是 Debug
使用 .cubin 文件代替 .ptx

如果有人仍然想就如何进一步提高性能提出一些想法 - 请分享他们！我乐于接受想法。不过这个问题已经解决了！

p.s。如果您的显示器以与 here 所述相同的方式闪烁，那么也请尝试增加延迟。

在 CUDA 内核中循环数据会导致应用程序中止

Looping over data in CUDA kernel causes app to abort

c#

loops

cuda

managed-cuda