выпуск:
По мере увеличения объема данных, обрабатываемых внутри цикла, который находится внутри CUDA kernel - это вызываетприложение, которое необходимо прервать!
исключение:
ManagedCuda.CudaException: 'ErrorLaunchFailed: на устройстве возникла исключительная ситуация при выполненииядро.Распространенные причины включают разыменование недействительного указателя устройства и доступ к разделяемой памяти за пределами.
вопрос:
Буду признателен, если кто-нибудьможет пролить свет на ограничения, которые я применяю в своей текущей реализации, и что именно вызывает сбой приложения ..
В качестве альтернативы, я прилагаю полный код ядра, если кто-то может сказать, как это можетбыть смоделированы таким образом, когда не выдается никаких исключений.Идея состоит в том, что ядро принимает 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, то код завершится успешно.
code - kernel (полная версия)
#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 данных в среде хоста.Поэтому я представляю, что независимо от того, сколько данных я копирую в графический процессор - оно не должно превышать 1Gb или даже максимум 11Gb.Позже я заметил, что сбой происходит только тогда, когда цикл внутри ядра имеет много записей данных для обработки.Это заставляет меня верить, что я нарушаю какие-то ограничения по времени ожидания потока или что-то в этом роде.Без надежного доказательства.
system
Мои системные характеристики 16Gb из Ram и GeForce 1080 Ti 11Gb.Используя Cuda 9.1. и managedCuda версию 8.0.22 (также пробовал с версией 9.x из главной ветки)
edit 1: 26.04.2018 Только что протестировал ту же логику, нотолько на OpenCL.Код не только успешно завершен, но и работает в 1,5-5 раз лучше, чем CUDA, в зависимости от размеров входных параметров:
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 / *Война 1095 * здесь.Если есть что-то, что меня должно волновать в моей первоначальной реализации CUDA - пожалуйста, дайте мне знать.
edit: 26.04.2018 .После следующих предложений из раздела комментариев мне удалось увеличить объем обрабатываемых данных до выдачи исключения в 3 раза.Я смог добиться этого, переключившись на .ptx, сгенерированный в режиме Release, а не Debug.Это улучшение может быть связано с тем, что в настройках Debug мы также установили Generate GPU Debug information на Yes и другие ненужные настройки, которые могут повлиять на производительность. Теперь я попытаюсь найти информацию о том, как можно увеличить время для ядра.. Я все еще не достигаю результатов OpenCL, но приближаюсь.
Для генерации файлов CUDA Я использую VS2017 Community, CUDA 9.1 project, v140 toolset, build for x64 платформа, события после сборки отключены, тип конфигурации: utility.Для генерации кода установлено: compute_30,sm_30.Я не уверен, почему это не sm_70, например.У меня нет других вариантов.