Почему фиктивная cudaMallo c ускоряет интервал в сборке?

Question

Я использую библиотеку moderngpu CUDA (https://github.com/moderngpu/moderngpu)

В moderngpu происходит странное ускорение работы функции interval_gather, если до cudaMalloc было выделено несколько байтов вызывая его.

Я вызываю метод mgpu::interval_gather(m0.data(), 1, m1.data(), 1, m2.data(), m3.data(), context); Каждый из m0, m1, m2, m3 - это mgpu::mem_t массив размером 1 000 000, заполненный 0 с.

Я делаю это 10 000 раз, и это займет около 3 секунд.

Однако, если я выделю некоторую фиктивную память перед ним:

cudaMalloc((void **)(&tmpPtr), sizeof(int));

Он ускоряется примерно в 10 раз, занимая 0,3 секунды. С чего бы это?

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

Однако, когда я попытался уменьшить размер массивов, я передаю в interval_gather, эффект прекращается. Затем это займет 0,3 секунды, независимо от того, был ли фиктивный алло c или нет.

Я запустил все это на GTX 980.

Вот полный фрагмент кода. Если для переменной magic установлено значение true, она будет ускоряться.

#include <iostream>
#include <chrono>
#include <string>

#include <moderngpu/kernel_intervalmove.hxx>

void print_ms(std::clock_t start, std::clock_t end, std::string desc)
{
    double ms = 1000.0 * (end - start) / CLOCKS_PER_SEC;
    std::cout << desc << ": " << ms << " ms." << std::endl;
}

void fun(bool magic, mgpu::context_t &context)
{
    int n = 10000000;
    mgpu::mem_t<int> m0(n, context);
    mgpu::mem_t<int> m1(n, context);
    mgpu::mem_t<int> m2(n, context);
    mgpu::mem_t<int> m3(n, context);

    void *tmpPtr;
    cudaMalloc((void **)(&tmpPtr), sizeof(int));
    if (!magic)
        cudaFree(tmpPtr);

    for (int aa = 0; aa < 10000; aa++)
        mgpu::interval_gather(m0.data(), 1,
                              m1.data(), 1,
                              m2.data(), 
                              m3.data(),
                              context);

    if (magic)
        cudaFree(tmpPtr);
}

int main(int argc, char *argv[])
{
    mgpu::standard_context_t context(false);
    std::clock_t c0 = std::clock();

    fun(false, context);
    context.synchronize();
    std::clock_t c1 = std::clock();
    print_ms(c0, c1, "1st");

    fun(false, context);
    context.synchronize();
    std::clock_t c2 = std::clock();
    print_ms(c1, c2, "2nd");

    fun(true, context);
    context.synchronize();
    std::clock_t c3 = std::clock();
    print_ms(c2, c3, "3rd");

    return 0;
}

Answer 1

Предисловие: я могу рассказать вам, что происходит, но не могу объяснить, почему.

TLDR: вызов cudaMalloc перед mgpu::interval_gather меняет производительность внутренних cudaMalloc и cudaFree звонки в пределах interval_gather operation примерно в 10 раз, и это разница в производительности, которую вы видите. Важным моментом является то, что память, которая выделяется перед вызовами mgpu , не должна освобождаться до тех пор, пока не будут выполнены вызовы .

Пример типичной синхронизации API одного вызова на mgpu::interval_gather в вашем код без волхвов c cudaMalloc вызов:

Трассировка API:

262.99ms  80.334us  cudaMalloc
263.07ms  11.912us  cudaLaunchKernel (_ZN4mgpu16launch_box_cta_kINS_15launch_params_tILi128ELi1ELi1ELi0EEENS_6detail11transform_fIS2_EEJZNS_21merge_path_partitionsILNS_8bounds_tE1ENS_19counting_iterator_tIiiEEPiNS_6less_tIiEEEENS_5mem_tIiEET0_lT1_llT2_RNS_9context_tEEUliE_mE
263.08ms  7.3310us  cudaLaunchKernel (_ZN4mgpu16launch_box_cta_kINS_12launch_box_tIJNS_7arch_20INS_12launch_cta_tILi128ELi11ELi9ELi0EEENS_7empty_tEEENS_7arch_35INS3_ILi128ELi7ELi5ELi0EEES5_EENS_7arch_52IS4_S5_EEEEEZNS_13transform_lbsIS5_ZNS_15interval_gatherIS5_PiSF_SF_SF_EE
263.09ms  51.357us  cudaFree

Трассировка GPU:

273.40ms  1.4400us              (1 1 1)       (128 1 1)        11        0B        0B  GeForce GTX 970         1         7  _ZN4mgpu16launch_box_cta_kINS_15launch_params_tILi128ELi1ELi1ELi0EEENS_6detail11transform_fIS2_EEJZNS_21merge_path_partitionsILNS_8bounds_tE1ENS_19counting_iterator_tIiiEEPiNS_6less_tIiEEEENS_5mem_tIiEET0_lT1_llT2_RNS_9context_tEEUliE_mEEEvSF_iDpSG_ [126]
273.40ms  5.7600us              (1 1 1)       (128 1 1)        56  5.5156KB        0B  GeForce GTX 970         1         7  _ZN4mgpu16launch_box_cta_kINS_12launch_box_tIJNS_7arch_20INS_12launch_cta_tILi128ELi11ELi9ELi0EEENS_7empty_tEEENS_7arch_35INS3_ILi128ELi7ELi5ELi0EEES5_EENS_7arch_52IS4_S5_EEEEEZNS_13transform_lbsIS5_ZNS_15interval_gatherIS5_PiSF_SF_SF_EEvT0_iT1_iT2_T3_RNS_9context_tEEUliiiNS_5tupleIJiEEESF_SF_E_SF_NSM_IJSF_EEEJSF_SF_EEEvSG_iSH_iSI_SL_DpSJ_EUliiSF_SF_E_JSF_SF_EEEvSG_iDpSH_ [127]

по сравнению с тем, когда волхвы c cudaMalloc используется:

Трассировка API:

2.32306s  4.7240us  cudaMalloc
2.32307s  7.9970us  cudaLaunchKernel (_ZN4mgpu16launch_box_cta_kINS_15launch_params_tILi128ELi1ELi1ELi0EEENS_6detail11transform_fIS2_EEJZNS_21merge_path_partitionsILNS_8bounds_tE1ENS_19counting_iterator_tIiiEEPiNS_6less_tIiEEEENS_5mem_tIiEET0_lT1_llT2_RNS_9context_tEEUliE_mE
2.32308s  6.7660us  cudaLaunchKernel (_ZN4mgpu16launch_box_cta_kINS_12launch_box_tIJNS_7arch_20INS_12launch_cta_tILi128ELi11ELi9ELi0EEENS_7empty_tEEENS_7arch_35INS3_ILi128ELi7ELi5ELi0EEES5_EENS_7arch_52IS4_S5_EEEEEZNS_13transform_lbsIS5_ZNS_15interval_gatherIS5_PiSF_SF_SF_EE
2.32308s  8.2070us  cudaFree

Трассировка графического процессора:

2.37275s  1.0240us              (1 1 1)       (128 1 1)        11        0B        0B  GeForce GTX 970         1         7  _ZN4mgpu16launch_box_cta_kINS_15launch_params_tILi128ELi1ELi1ELi0EEENS_6detail11transform_fIS2_EEJZNS_21merge_path_partitionsILNS_8bounds_tE1ENS_19counting_iterator_tIiiEEPiNS_6less_tIiEEEENS_5mem_tIiEET0_lT1_llT2_RNS_9context_tEEUliE_mEEEvSF_iDpSG_ [120129]
2.37276s  4.0000us              (1 1 1)       (128 1 1)        56  5.5156KB        0B  GeForce GTX 970         1         7  _ZN4mgpu16launch_box_cta_kINS_12launch_box_tIJNS_7arch_20INS_12launch_cta_tILi128ELi11ELi9ELi0EEENS_7empty_tEEENS_7arch_35INS3_ILi128ELi7ELi5ELi0EEES5_EENS_7arch_52IS4_S5_EEEEEZNS_13transform_lbsIS5_ZNS_15interval_gatherIS5_PiSF_SF_SF_EEvT0_iT1_iT2_T3_RNS_9context_tEEUliiiNS_5tupleIJiEEESF_SF_E_SF_NSM_IJSF_EEEJSF_SF_EEEvSG_iSH_iSI_SL_DpSJ_EUliiSF_SF_E_JSF_SF_EEEvSG_iDpSH_ [120130]

Вы можете ясно видеть, что производительность cudaMalloc и cudaFree сильно меняется, но не более того:

Интересно, что если вы оставите выделенную память magi c, изменение производительности будет сохраняться между вызовами вашей тестовой функции, например:

void fun(bool magic, mgpu::context_t &context)
{
    int n = 10000000;
    mgpu::mem_t<int> m0(n, context);
    mgpu::mem_t<int> m1(n, context);
    mgpu::mem_t<int> m2(n, context);
    mgpu::mem_t<int> m3(n, context);

    void *tmpPtr = 0;
    if (magic) cudaMalloc((void **)(&tmpPtr), sizeof(int));

    for (int aa = 0; aa < 10000; aa++)
        mgpu::interval_gather(m0.data(), 1,
                              m1.data(), 1,
                              m2.data(), 
                              m3.data(),
                              context);

}

int main(int argc, char *argv[])
{
    {
    mgpu::standard_context_t context(false);
    std::clock_t c0 = std::clock();

    fun(false, context);
    context.synchronize();
    std::clock_t c1 = std::clock();
    print_ms(c0, c1, "1st");

    fun(true, context);
    context.synchronize();
    std::clock_t c2 = std::clock();
    print_ms(c1, c2, "2nd");

    fun(false, context);
    context.synchronize();
    std::clock_t c3 = std::clock();
    print_ms(c2, c3, "3rd");
    }

    return 0;
}

делает это:

$ nvcc -arch=sm_52 -std=c++11 --expt-extended-lambda -I ~/mgpu/moderngpu/src -o mgpuspeed mgpuspeed.cu 
$ ./mgpuspeed 
1st: 1287.37 ms.
2nd: 201.205 ms.
3rd: 202.275 ms.

[Все тайминги на GTX970, CUDA 10.1, Ubuntu 18.04 LTS, драйвер 440.59]

Это несмотря на то, что каждый mem_t конструктор call вызывает cudaMalloc call и соответствующий cudaFree call, когда объекты выпадают из области видимости. Так что это что-то, относящееся в отдельности к маленькому вызову "волхвов c" cudaMalloc, а не к cudaMalloc в целом. Это заставляет меня думать, что изменение производительности на самом деле происходит в диспетчере памяти, а не является каким-то тонким изменением в планировании или асинхронном потоке программы. Это может быть что-то столь же простое, как перемещение распределителя slab-файлов на другое выравнивание или размер страницы, или что-то, что изменяет его работу для внутренних вызовов cudaMalloc.

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