Память CUDA (тип) для вычисления только устройства во время вызовов ядра (вычисление 1.1 или 1.2)

Question

В настоящее время я изучаю CUDA, и мой алгоритм должен выполнять некоторые тяжелые вычисления на основе некоторых входных данных. Эти вычисления производятся в цикле, который вращается до 1024 раундов. Все работает нормально, пока у меня есть небольшое количество потоков (< 100'000) на каждое ядро, но если я хочу использовать больше потоков, ядро ​​будет прервано окнами, так как требуется слишком много времени для завершения.

Мое решение было бы разделить тяжелые расчеты в нескольких вызовов ядра:

1. главный ядро, которое подготавливает входные данные и вычисляет первые х раундов (петля раскатали). Это будет вызываться только один раз на каждый вход.
2. работа Ядро, которое выполняет следующие x раундов (цикл развернут). Это будет вызываться так часто, как требуется для расчета всех необходимых раундов.

Между каждым вызов ядра (один главным, многие работы), я должен сохранить 16 + байты данных, которые будут использоваться при следующем вызове (длина длина ввода, он фиксируется на главный звонок). Основное ядро ​​ядра начнет записывать эти байты, а ядро ​​ будет читать их, запускать следующие вычисления и записывать исходные данные с новым результатом. Мне нужны только эти данные на устройстве, не требуется хост-доступ. Какую память я должен использовать для этого? По крайней мере, это должна быть глобальная память, поскольку она является единственной записываемой памятью, которая сохраняется во время вызовов ядра, не так ли? Но тогда, что? Не могли бы вы дать мне совет о том, как мне приступить к работе с правильной памятью (и лучшая производительность)?

В «псевдокод», это может выглядеть следующим образом:

prepare memory to hold threads * (16 + length) bytes 

for length = 1 to x step 1 
    call mainKernel 
    rounds = 1024 - rounds_done_in_main 
    for rounds to 0 step rounds_done_in_work 
    call workKernel 
    end for 
end for 

cleanup memory 

-------- 

template <unsigned char length> __global__ mainKernel() { 
    unsigned char input[length]; 
    unsigned char output[16]; 
    const int tid = ...; 

    devPrepareInput<length>(input); 

    calc round 1: doSomething<length>(output, input) 
    calc round 2: doSomething<length>(output, output + input) // '+' == append 

    write data to memory based on tid // data == output + input 
} 

template <unsigned char length, remaining rounds> __global__ workKernel() { 
    unsigned char *input; 
    unsigned char *output; 
    const int tid = ...; 

    read data from memory based on tid 
    ouput = data 
    input = data+16 

    if rounds >= 1 
    calc round x : doSomething<length>(output, output + input) 
    if rounds >= 2 
    calc round x+1: doSomething<length>(output, output + input) // '+' == append 

    if rounds == x // x is the number of rounds in the last work call 
    do final steps on output 
    else 
    write ouput + input to memory based on tid (for next call) 
} 

Answer 1

Да, вы можете сделать это с памятью устройства. Переменная, объявленная с __device__, предоставляет статическое объявление буфера, которое может использоваться непосредственно ядрами, без каких-либо операций cudaMemcpy, и не требует, чтобы указатель был явно передан ядру. Поскольку он имеет номер lifetime of the application, данные в нем будут сохраняться с одного вызова ядра на следующий.

#define NUM_THREADS 1024 
#define DATA_PER_THREAD 16 
__device__ int temp_data[NUM_THREADS*DATA_PER_THREAD]; 

__global__ my_kernel1(...){ 
    int my_data[DATA_PER_THREAD] = {0}; 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    // perform calculations 

    // write out temp data 
    for (int i = 0; i < DATA_PER_THREAD; i++) temp_data[i + (idx * DATA_PER_THREAD)] = my_data[i]; 
    } 

__global__ my_kernel2(...){ 
    int my_data[DATA_PER_THREAD]; 
    // read in temp data 
    for (int i = 0; i < DATA_PER_THREAD; i++) my_data[i] = temp_data[i + (idx * DATA_PER_THREAD)]; 
    // perform calculations 

    } 

Существует множество способов, которыми вы могли бы оптимизировать это на основе шаблона использования в ядре. Передача данных в/из my_data действительно не требуется. Очевидно, ваш код ядра мог бы просто получить доступ к temp_data прямо вместо my_data с соответствующей индексацией.

Если вы решили загрузить/сохранить его, вы можете чередовать данные, чтобы разрешить совместный доступ во время чтения и записи данных for.