Проблемы с потоками CUDA

Question

Я запускаю CUBLAS v2.0 в разных потоках на одном графическом процессоре (Tesla C2050) путем разделения входных матриц (A [x / num_of_streams * y] B [x y] = C [x / num_of_streams * y]), но как-то это занимает больше времени, когда я использую потоки CUDA. Вот фрагмент кода:

             //plan is a struct containing the matrix dimensions and stream numbers
             //parallel in nstreams - should be! MAX 16 streams could run concurrently
            //Copy A - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyA_in_streams (&plan[i]);
            //Copy B - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyB_in_streams (&plan[i]);

            //Create handles - serial
            for(i = 0; i < nstreams; i++)
                    handle[i] = create_handle();

            //Run kernels - first doing a cublasSetStream(handle, plan->stream) before running cublasDgemm... 
            for(i = 0; i < nstreams; i++)
                    cudgemm_kernel_in_streams (&plan[i], handle[i], 1.0f, 1.0f);

            //Destroy handles - serial
            for(i = 0; i < nstreams; i++)
                    destroy_handle (handle[i]);

            //Copy C - cudaMemCpyAsync
            for(i = 0; i < nstreams; i++)
                    cudgemm_copyC_in_streams (&plan[i]);

            //EDIT: Function body

            //The other two copy functions are exactly the same as this
            void cudgemm_copyA_in_streams(TGPUplan *plan)
           {
                 cudasafe(cudaMemcpyAsync(plan->Ad_Data, plan->Ah_Data, (plan->Acols * plan->Arows * sizeof(double)), cudaMemcpyHostToDevice, plan->stream) );

            }

            //Create handle
            cublasHandle_t create_handle ()
            {
                   cublasHandle_t handle;
                   checkError(cublasCreate(&handle), "cublasCreate() error!\n");
                   return handle;
             }

             //Destroy handle
             void destroy_handle (cublasHandle_t handle)
             {
                  checkError(cublasDestroy(handle), "cublasDestroy() error!\n");
             }

             //Kernel
             void cudgemm_kernel_in_streams(TGPUplan *plan, cublasHandle_t handle, const double alpha, const double beta)
             {
                   cublasStatus_t ret;
                   cublasSetStream(handle, plan->stream);

                   ret = cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, plan->Arows, plan->Ccols, plan->Acols, &alpha, plan->Ad_Data, plan->Arows, plan->Bd_Data, plan->Brows, &beta, plan->Cd_Data, plan->Crows);
                   checkError(ret, "cublas Dgemm returned an error!\n");
              }

Таким образом, я отскакиваю назад и вперед между потоками и назначаю работу, ожидая лучшего времени выполнения, но я замечаю, что чем больше количество потоков, тем больше времени занимает программа по сравнению с версией, которая не использует поток. Куда я иду не так? Перекрестное сообщение на форумах Nvidia - http://forums.nvidia.com/index.php?showtopic=209420

EDIT:

Я изменил свою программу следующим образом:

            //copy data
            for(i = 0; i < nstreams; i++)
            {
                    cudgemm_copyA_in_streams (&plan[i]);
                    cudgemm_copyB_in_streams (&plan[i]);
            }

            //Run kernel and copy back
            for(i = 0; i < nstreams; i++)
            {
                    cudgemm_kernel_in_streams (&plan[i], handle[i], 1.0f, 1.0f);
                    cudgemm_copyC_in_streams (&plan[i]);
            }

Когда я профилирую свою программу для порядка матрицы 6144, я получаю следующую информацию:

Kernel time = 42.75 % of total GPU time 
Memory copy time = 28.9 % of total GPU time
Kernel taking maximum time = fermiDgemm_v2_kernel_val (42.8% of total GPU time)
Memory copy taking maximum time = memcpyHtoDasync (21.7% of total GPU time)
Total overlap time in GPU = 65268.3 micro sec. (3.6% of total GPU time)

Когда я проверяю вышеуказанный цикл, я получаю время 0,000284 с против 1,703289 с для версии, которая не использует потоки (в этой версии я также синхронизирую две последовательные копии памяти, вызов ядра и оставшийся memCpy) , Я думаю, что поскольку я не использую какие-либо конструкции синхронизации, возможно, я печатаю время до фактического завершения вычислений (мне трудно поверить, что улучшение на 100%).

Answer 1

Я предлагаю два изменения:

1) перенести создание / уничтожение дескриптора cuBLAS за пределы копий и вызовов ядра. Возможно, он нарушает параллелизм, выполняя ненужную синхронизацию контекста.

2) сделать memcpy вместе в одном цикле над потоками. Таким образом, копия B потока 0 не выполняет никакой дополнительной синхронизации, чтобы дождаться завершения A memcpy. т.е. сделать это:

        for(i = 0; i < nstreams; i++) {
                cudgemm_copyA_in_streams (&plan[i]);
                cudgemm_copyB_in_streams (&plan[i]);
        }

не это:

        for(i = 0; i < nstreams; i++)
                cudgemm_copyA_in_streams (&plan[i]);
        for(i = 0; i < nstreams; i++)
                cudgemm_copyB_in_streams (&plan[i]);

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

Answer 2

Я бы также предложил проверить РАЗМЕР копий, вы должны начать использовать разные потоки только тогда, когда время для передачи одного блока памяти можно сравнить со временем, необходимым для его вычисления.Если время для передачи мало по сравнению со временем вычисления, то добавление потоков добавляет больше накладных расходов при их управлении.Используйте визуальный профилировщик, чтобы узнать, сколько времени занимает выполнение различных шагов.Сделайте график с различными входами памяти.