V_oron

CUDA: blocks synchronisation

Sep 20th, 2012
558
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  4. #include <cstdlib>
  5.  
  6. const int REPEAT_TEST           = 10;
  7. const int TRANSF_ITER           = 100;
  8. const int BLOCKS []             = {  2,   4,   8, 16, 32};                  // проверки на соразмерность масивов нет!
  9. const int THREADS_PER_BLOCK []  = {512, 256, 128, 64, 32};
  10. const int THREADS_PER_BLOCK_max = 512;
  11.  
  12.  
  13. __device__ volatile unsigned int count; //volatile
  14. __device__ void do_while_count_not_eq(int val)
  15. {
  16.     /*
  17.     // Цикл задержки в виде ассемблерной вставки. Точно будет в коде, не требует "volatile"-переменной count.
  18.     asm("{\n\t"
  19.         "$my_while_label: \n\t"
  20.         "   .reg .u32      r_count;            \n\t"
  21.         "   .reg .pred     p;                  \n\t"
  22.         "   ld.global.u32  r_count, [count];   \n\t"
  23.         "   setp.ne.u32    p, r_count, %0;     \n\t"
  24.         "@p bra            $my_while_label;    \n\t"
  25.         "}\n\t"
  26.         : : "r"(val));
  27.     */
  28.    
  29.     // Цикл задержки просто на Си. Компилятор nvcc "выкинет" его, если переменная count не определена с использованием "volatile".
  30.     while (count != val);
  31. }
  32. /* Фунция начальной инициализации флага-счетчика: */
  33. __device__ void InitSyncWholeDevice(const int index)
  34. {
  35.     if (index == 0)                                                         // Первый поток в grid`е запишет нулем (начальным значением) флаг-счетчик блоков.
  36.         count = 0;
  37.        
  38.     if (threadIdx.x == 0)                                                   // Первый поток каждого block`а будет ждать, пока флаг-счетчик дейcтвительно станет нулем.
  39.         do_while_count_not_eq(0);
  40.        
  41.     __syncthreads();                                                        // Заставляем остальные потоки каждого block`а ждать, пока первые не выйдут из цикла.
  42. }                                                                           // Все, флаг-аккумулятор записан. Все потоки на device более-менее идут вровень.
  43. /* Фунция синхронизации потоков на device: */
  44. __device__ void SyncWholeDevice()
  45. {
  46.     unsigned int oldc;                                                      // Переменная под значение счетчика до инкремента.
  47.     __threadfence();                                                        // Каждый поток ждет, пока записанное им в gmem и smem, станет видно всему grid`у. __syncthreads отвечает только за видимость для block.
  48.     if (threadIdx.x == 0)                                                   // Первые потоки каждого block`а атомарным образом инкрементируют (каждый по разу) флаг-аккумулятор.
  49.     {
  50.         oldc = atomicInc((unsigned int *)&count, gridDim.x-1);              // В oldc кладется значение count до "+1".
  51.         __threadfence();                                                    // Пусть поток подождет, пока его инкремент "дойдет" до ячейки в gmem.
  52.         if (oldc != (gridDim.x-1))                                          // Если это последний блок (остальные уже инкрементировали count и ждут за счет цикла ниже), то и незачем ему считывать count. Так как предварительно убедились, что его инкремент записан в gmem.
  53.             do_while_count_not_eq(0);                                       // Каждый первый поток в каждом block`е зациклен, пока все блоки не закончат работу.
  54.     }
  55.     __syncthreads();                                                        // Заставляем потоки в каждом блоке ждать, пока первые не выйдут из цикла.
  56. }
  57.  
  58.  
  59. __device__ void SingleTransform(const float *src, float *dest)
  60. {
  61.     float a = 0;
  62.  
  63.     /* Цикл вычисления элемента dest[]. В двух вариантах. */
  64.     // Вариант №1: без использования разделяемой памяти:
  65.         for(int j = 0; j < gridDim.x*blockDim.x; ++j)
  66.             a += src[j];
  67.     // Вариант №2: с использованием разделяемой памяти:
  68.         /*__shared__ float smemBuff [THREADS_PER_BLOCK_max];
  69.         for(int jb = 0; jb < gridDim.x; ++jb)
  70.         {
  71.             __syncthreads();                                                // Необходимо дождться, пока smemBuff[] больше не будет нужен никакому потоку из блока.
  72.             smemBuff[threadIdx.x] = src[threadIdx.x + jb*blockDim.x];
  73.             __syncthreads();                                                // Необходимо дождться, пока в smemBuff[] все потоки блока запишут значения.
  74.             for(int jt = 0; jt < blockDim.x; ++jt)
  75.                 a += smemBuff[jt];
  76.         }*/
  77.  
  78.     a /= (gridDim.x*blockDim.x);
  79.     dest[threadIdx.x + blockIdx.x*blockDim.x] = a;
  80. }
  81.  
  82. __global__ void Kernel_SKL(float *X, float *P)
  83. {
  84.     InitSyncWholeDevice(threadIdx.x + blockIdx.x*blockDim.x);
  85.     for (int i = 0; i < TRANSF_ITER; ++i)
  86.     {
  87.         SingleTransform(X, P);
  88.         SyncWholeDevice();
  89.         SingleTransform(P, X);
  90.         SyncWholeDevice();
  91.     }
  92. }
  93.  
  94. __global__ void Kernel_MKL(float *src, float *dest)
  95. {
  96.     SingleTransform(src, dest);
  97. }
  98.  
  99.  
  100. int main(void)
  101. {
  102.     int NUM_OF_ELEMS;
  103.     float *X_h;
  104.     float *X_d, *P_d;
  105.     float *res[REPEAT_TEST+1];
  106.  
  107.     float elapsedTime_MKL[2];
  108.     float elapsedTime_SKL[2];
  109.        
  110.     cudaEvent_t start, stop;
  111.     cudaEventCreate(&start);
  112.     cudaEventCreate(&stop);
  113.    
  114.     FILE *fo = fopen("out.txt", "w");
  115.    
  116.     for(int jnb = 0; jnb < sizeof(BLOCKS)/sizeof(*BLOCKS); ++jnb)
  117.     for(int jnt = 0; jnt < sizeof(THREADS_PER_BLOCK)/sizeof(*THREADS_PER_BLOCK); ++jnt)
  118.     {
  119.         printf("BLOCKS = %3i, THREADS_PER_BLOCK = %3i:\n", BLOCKS[jnb], THREADS_PER_BLOCK[jnt]);
  120.         NUM_OF_ELEMS = BLOCKS[jnb] * THREADS_PER_BLOCK[jnt];
  121.        
  122.         X_h = new float [NUM_OF_ELEMS];
  123.         cudaMalloc(&X_d, NUM_OF_ELEMS*sizeof(float));
  124.         cudaMalloc(&P_d, NUM_OF_ELEMS*sizeof(float));
  125.        
  126.         srand(time(NULL));
  127.         for(int j = 0; j < NUM_OF_ELEMS; ++j)
  128.             X_h[j] = float(rand()) / RAND_MAX; // как вариант: float(j+1)/NUM_OF_ELEMS;
  129.        
  130.         for(int r = 0; r <= REPEAT_TEST; ++r)
  131.             res[r] = new float [NUM_OF_ELEMS];
  132.        
  133.         elapsedTime_MKL[0] = elapsedTime_SKL[0] = 0;
  134.         for(int r = 1; r <= REPEAT_TEST; ++r)
  135.         {
  136.             printf("  r = %i/%i\r", r, REPEAT_TEST);
  137.             fflush(stdout);
  138.            
  139.             cudaMemcpy(X_d, X_h, NUM_OF_ELEMS*sizeof(float), cudaMemcpyHostToDevice);
  140.             cudaEventRecord(start, 0);
  141.             for (int i = 0; i < TRANSF_ITER; ++i)
  142.             {
  143.                 Kernel_MKL <<<BLOCKS[jnb], THREADS_PER_BLOCK[jnt]>>> (X_d, P_d);
  144.                 Kernel_MKL <<<BLOCKS[jnb], THREADS_PER_BLOCK[jnt]>>> (P_d, X_d);
  145.             }
  146.             cudaEventRecord(stop, 0);
  147.             cudaEventSynchronize(stop);
  148.             cudaEventElapsedTime(&elapsedTime_MKL[1], start, stop);
  149.             elapsedTime_MKL[0] += elapsedTime_MKL[1];
  150.             cudaMemcpy(res[0], X_d, NUM_OF_ELEMS*sizeof(float), cudaMemcpyDeviceToHost);    // Дабы не усложнять пример, лишние несколько раз скопируем одно и то же в res[0].
  151.    
  152.             cudaMemcpy(X_d, X_h, NUM_OF_ELEMS*sizeof(float), cudaMemcpyHostToDevice);
  153.             cudaEventRecord(start, 0);
  154.             Kernel_SKL <<<BLOCKS[jnb], THREADS_PER_BLOCK[jnt]>>> (X_d, P_d);
  155.             cudaEventRecord(stop, 0);
  156.             cudaEventSynchronize(stop);
  157.             cudaEventElapsedTime(&elapsedTime_SKL[1], start, stop);
  158.             elapsedTime_SKL[0] += elapsedTime_SKL[1];
  159.             cudaMemcpy(res[r], X_d, NUM_OF_ELEMS*sizeof(float), cudaMemcpyDeviceToHost);
  160.         }
  161.         printf("  Elapsed Time\n");
  162.         printf("    multi kerenel launch:  %20.14f\n", elapsedTime_MKL[0]/REPEAT_TEST);
  163.         printf("    single kerenel launch: %20.14f\n", elapsedTime_SKL[0]/REPEAT_TEST);
  164.         fprintf(fo, "%i  %i  %20.14f   %20.14f\n", jnb, jnt, elapsedTime_MKL[0]/REPEAT_TEST, elapsedTime_SKL[0]/REPEAT_TEST);
  165.        
  166.         // Сравним между собой все j-е элементы из каждой r-ой реализации. Выявит численную ошибку за счет использования неактуальных значений элементов массивов.
  167.         for(int j = 0; j < NUM_OF_ELEMS; ++j)
  168.         {
  169.             int r;
  170.             float e = res[0][j];
  171.             for(r = 1; r < REPEAT_TEST; ++r)
  172.                 if (e != res[r][j])
  173.                 {
  174.                     printf("    Numeric error of SKL-method has been detected!\n");
  175.                     break;
  176.                 }
  177.             if (e != res[r][j])
  178.                 break;
  179.         }
  180.        
  181.         cudaFree(X_d);
  182.         cudaFree(P_d);
  183.         delete[] X_h;
  184.         for(int r = 0; r < (REPEAT_TEST+1); ++r)
  185.             delete[] res[r];
  186.     }
  187.    
  188.     fclose(fo);
  189.    
  190.     return 0;
  191. }
RAW Paste Data