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

问题：

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

例外：

ManagedCuda.CudaException：’ErrorLaunchFailed：执行内核时设备发生exception。 常见原因包括解除引用无效设备指针和访问超出范围的共享内存。

题：

如果有人可以了解我正在使用当前实现的限制以及究竟是什么导致应用程序崩溃，我将不胜感激。

或者，我附加了一个完整的内核代码，以便有人可以说如何在没有抛出exception的情况下以这种方式重新建模。 我们的想法是，内核正在接受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，那么应用程序将开始中止。 使用相同类型的输入参数，如果我们将arraySize覆盖或硬核为10_000，则代码将成功完成。

代码 – 内核（完整版）

 #iclude  #include "cuda_runtime.h" #include  #include  #include  #define _SIZE_T_DEFINED #ifndef __CUDACC__ #define __CUDACC__ #endif #ifndef __cplusplus #define __cplusplus #endif texture 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 = 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]; // ie bool var out_dev = new CudaDeviceVariable(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]; // ie bool var in1_dev = new CudaDeviceVariable(in1_host.Length); var in2_dev = new CudaDeviceVariable(in2_host.Length); var in3_dev = new CudaDeviceVariable(in3_host.Length); var in4_dev = new CudaDeviceVariable(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 。 后来我注意到只有当内核中的循环有许多要处理的数据记录时才会发生崩溃。 这让我相信我遇到某种线程超时限制或类似的东西。 没有坚实的证据。

系统

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

编辑1：26.04.2018刚刚测试了相同的逻辑，但仅限于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 = 0) { output[index] = true; } } 

我真的不想在这里开始OpenCL / CUDA战争。 如果在我原来的CUDA实施中有任何我应该关注的事情 – 请告诉我。

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

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

Answer 1

我已经设法提高了OpenCL的CUDA性能。 而且更重要的是 – 代码现在可以完成执行而无需例外。 学分归Robert Crovella所有 。 谢谢！

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

• CPU Intel i7 8700k 12核（6 + 6）
• GPU GeForce 1080 Ti 11Gb

这是我的结果（图书馆/技术）：

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

解决方案1：

解决方案是将WDDM TDR Delay从默认的2秒增加到10秒。 就这么简单 。

解决方案2：

我能够通过以下方式挤出更多的表现：

1. 在CUDA项目属性中将compute_30,sm_30设置更新为compute_30,sm_30

2. 使用Release设置而不是Debug

3. 使用.cubin文件而不是.ptx

如果有人还想提出一些关于如何进一步提高性能的想法 – 请分享！ 我很开心。 不过，这个问题已经解决了！

ps如果您的显示器以与此处所述相同的方式闪烁，则尝试增加延迟。