questão:

À medida que aumenta a quantidade de dados que estão sendo processados dentro do loop que está dentro doCUDA kernel - faz com que o aplicativo seja abortado!

exceção:

ManagedCuda.CudaException: 'ErrorLaunchFailed: Ocorreu uma exceção no dispositivo durante a execução de um kernel. As causas comuns incluem desreferenciar um ponteiro de dispositivo inválido e acessar a memória compartilhada fora dos limites.

Pergunta, questão:

Gostaria que alguém pudesse esclarecer as limitações que estou atingindo com minha implementação atual e o que exatamente causa o travamento do aplicativo.

Como alternativa, estou anexando um código do kernel completo, se alguém puder dizer como ele pode ser re-modelado de tal maneira, quando nenhuma exceção for lançada. A idéia é que o kernel esteja aceitandocombinations e, em seguida, executando cálculos no mesmo conjunto dedata (em um loop). Portanto, os cálculos de loop que estão dentro devem ser seqüenciais. A sequência na qual o próprio kernel é executado é irrelevante. É problema combinatório.

Qualquer conselho é bem-vindo.

código (versão curta, suficiente para abortar o aplicativo):

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.
        }
    }
}

No exemplo acima, se oarraySize estiver próximo de 50_000, o aplicativo será interrompido. Com o mesmo tipo de parâmetros de entrada, se substituirmos ouRadical aarraySize para 10_000, o código termina com êxito.

code - kernel (versão completa)

#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;
        }
    }
}

código - host (versão completa)

    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);
    }

análise

Minha suposição inicial foi que eu estou tendo problemas de memória, no entanto, de acordo com o depurador do VS, estou atingindo um pouco acima500mb de dados em um ambiente host. Então, imagino que, não importa quantos dados eu copie para a GPU - ele não deve exceder1Gb ou até no máximo11Gb. Mais tarde, notei que a falha ocorre apenas quando o loop que está dentro de um kernel está tendo muitos registros de dados para processar. Isso me faz acreditar que estou atingindo algum tipo de limite de tempo limite de thread ou algo desse tipo. Sem uma prova sólida.

sistema

As especificações do meu sistema são16Gb doRameGeForce 1080 Ti 11Gb. UsandoCuda 9.1.emanagedCuda versão8.0.22 (também tentei com a versão 9.x da ramificação principal)

edit 1: 26.04.2018 Acabei de testar a mesma lógica, mas apenas emOpenCL. O código não apenas foi concluído com êxito, mas também apresenta um tempo de 1,5 a 5 vezes melhor que oCUDA, dependendo dos tamanhos dos parâmetros de entrada:

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;
    }
}

Eu realmente não quero começarOpenCL/CUDA guerra aqui. Se houver algo que eu deveria me preocupar no meu originalCUDA implementação - por favor me avise.

editar: 26.04.2018. Depois de seguir as sugestões da seção de comentários, fui capaz de aumentar a quantidade de dados processados, antes que uma exceção fosse lançada, em 3x. Consegui isso mudando para.ptx gerado emRelease modo, em vez deDebug modo. Essa melhoria pode estar relacionada ao fato de que, emDebug configurações também temosGenerate GPU Debug information definido comoYes e outras configurações desnecessárias que podem afetar o desempenho. Agora, tentarei pesquisar informações sobre como os tempos podem ser aumentados para o kernel. Ainda não estou alcançando os resultados deOpenCL, mas chegando perto.

ParaCUDA geração de arquivo que estou usandoVS2017 Community, CUDA 9.1 projeto,v140 toolsetconstruir parax64 plataforma, eventos pós-compilação desativados, tipo de configuração:utility. Geração de código definida como:compute_30,sm_30. Não sei por que não ésm_70, por exemplo. Eu não tenho outras opções.