Преимущества разделения большого ядра CUDA и использования динамического параллелизма

Question

У меня есть большое ядро, в котором начальное состояние развивается с использованием разных методов. То есть, у меня есть цикл в ядре, в этом цикле определенный предикат оценивается по текущему состоянию и по результату этого предиката выполняется определенное действие.

Ядро требует немного временных данных и разделяемой памяти, но поскольку он большой, он использует 63 регистра, а заполнение очень мало.

Я хотел бы разделить ядро ​​во многих маленьких ядрах, но каждый блок полностью независим от остальных, и я (думаю, я) не могу использовать один поток в главном коде для запуска нескольких небольших ядер.

Я не уверен, что потоки адекватны для такого рода работ, я никогда не использовал их, но поскольку у меня есть возможность использовать динамический параллелизм, я бы хотел, если бы это был хороший вариант для реализации такого рода работы , Быстро ли запустить ядро ​​из ядра? Нужно ли копировать данные в глобальной памяти, чтобы сделать их доступными для подъязыка?

Если я разбил свое большое ядро ​​во многих маленьких и оставил первое ядро ​​с основным циклом, который при необходимости вызвал необходимое ядро ​​(что позволяет мне перемещать временные переменные в каждом подъярусе), поможет мне увеличить размещение?

Я знаю, что это немного общий вопрос, но я не знаю эту технологию, и я бы хотел, чтобы она соответствовала моему делу или потокам было лучше.

EDIT: Чтобы предоставить некоторые другие детали, вы можете себе представить мое ядро, чтобы иметь такую ​​структуру:

__global__ void kernel(int *sampleData, int *initialData) { 
    __shared__ int systemState[N]; 
    __shared__ int someTemp[N * 3]; 
    __shared__ int time; 
    int tid = ...; 
    systemState[tid] = initialData[tid]; 

    while (time < TIME_END) { 
     bool c = calc_something(systemState); 
     if (c) 
      break; 
     someTemp[tid] = do_something(systemState); 
     c = do_check(someTemp); 
     if (__syncthreads_or(c)) 
      break; 
     sample(sampleData, systemState); 
     if (__syncthreads_and(...)) { 
      do_something(systemState); 
      sync(); 
      time += some_increment(systemState); 
     } 
     else { 
      calcNewTemp(someTemp, systemState); 
      sync(); 
      do_something_else(someTemp, systemState); 
      time += some_other_increment(someTemp, systemState); 
     } 
    } 
    do_some_stats(); 
} 

это, чтобы показать вам, что есть основной цикл, что есть временные данные, которые используются где-то, а не в других точках, что есть общие данные, точки синхронизации и т. д.

Темы используются для вычисления векторных данных, хотя в каждом блоке имеется, в идеале, одна петля (ну, конечно, это это не правда, но логически это так) ... Один «большой поток» для каждого блока.

Теперь я не уверен, как использовать потоки в этом случае ... Где находится «большая петля»? На хосте, я думаю ... Но как мне скопировать из одного цикла все блоки? Это то, что оставляет меня самым сомнительным. Могу ли я использовать потоки из разных потоков хоста (один поток на блок)?

Я менее сомнительно отношусь к динамическому параллелизму, потому что я мог бы легко поддерживать большую петлю, но я не уверен, могу ли я иметь преимущества здесь.

Answer 1

Я извлек выгоду из динамического параллелизма для решения задачи интерполяции вида:

int i = threadIdx.x + blockDim.x * blockIdx.x; 

for(int m=0; m<(2*K+1); m++) { 

    PP1 = calculate_PP1(i,m); 
    phi_cap1 = calculate_phi_cap1(i,m); 

     for(int n=0; n<(2*K+1); n++) { 

      PP2 = calculate_PP2(i,m); 
      phi_cap2 = calculate_phi_cap2(i,n); 

      atomicAdd(&result[PP1][PP2],data[i]*phi_cap1*phi_cap2); } } } 

где K=6. В этой задаче интерполяции вычисление каждого сложения не зависит от других, поэтому я разделил их на ядро ​​(2K+1)x(2K+1).

Из моего (возможно, неполного) опыта динамический параллелизм поможет, если у вас есть несколько независимых итераций. Для большего количества итераций, возможно, вы можете в конечном итоге вызвать дочернее ядро ​​несколько раз, и поэтому вы должны проверить, будет ли накладные расходы при запуске ядра ограничивающим фактором.