Я пытаюсь использовать динамический параллелизм для улучшения алгоритма, который я имею в CUDA. В моем исходном решении CUDA каждый поток вычисляет число, общее для каждого блока. Я хочу сначала запустить грубое (или с низким разрешением) ядро, где потоки вычисляют общее значение только один раз (например, если каждый поток представляет один блок). Затем каждый поток создает небольшую сетку из 1 блока (потоки 16x16) и запускает для нее дочернее ядро, передавая общее значение. Теоретически это должно быть быстрее, потому что каждый сохраняет много избыточных операций. Но на практике решение работает очень медленно, я не знаю почему.

Это очень упрощенный код, просто идея.

__global__ coarse_kernel( parameters ){
    int common_val = compute_common_val();
    dim3 dimblock(16, 16, 1);
    dim3 dimgrid(1, 1, 1);
    child_kernel <<< dimgrid, dimblock >>> (common_val, parameters);

}

__global__ child_kernel( int common_val, parameters ){
    // use common value
    do_computations(common_val, parameters);
}

Количество child_kernels много, по одному на поток, и должно быть около 400x400 потоков. Из того, что я понимаю, GPU должен обрабатывать все эти ядра параллельно, верно?

Или дочерние ядра обрабатываются как-то последовательно?

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