Комментировать

CUDA search

Всем доброго времени суток!
Почти через месяц защита диссертации. Разбираться во всех дебрях программирования на CUDA уже не осталось времени. Нужна программа, реализующая алгоритм поиска на GPU. Какой именно алгоритм это тоже большой вопрос. Алгоритмов поиска не мало и какой из них лучше распараллелиться на GPU?
Может кто подскажет как реализовать? Буду очень рад любой помощи. Спасибо.

Пока набросал простой линейный поиск. замечаниям и отзывам тоже буду рад.

При создании программы использовались следующие средства и инструменты:
1. Видеокарта из серии nVidia GeForce 8xxx/9xxx или более современная.
2. CUDA Toolkit v.2.3 (скачать можно здесь: www.nvidia.ru/object/cuda_get_ru.html)
3. CUDA SDK v.2.3 (скачать можно там же где Toolkit)
4. Visual Studio 2008
5. CUDA Visual Studio Wizard (скачать можно здесь: sourceforge.net/projects/cudavswizard/)

Текст программы

  1. #include <stdio.h>
  2. #include <time.h>
  3. #define CPU 1
  4. #define GPU 2
  5. // функция выполняемая на центральном процессоре
  6. __host__ void searchCPU ( float * data, int * result )
  7. {
  8.         int len = 16 * 1024 * 1024;    
  9.  
  10.         for ( int j = 0; j < len; j++ )
  11.         {
  12.                 if (data[j] == 2.0f )
  13.                 {
  14.                         result[0] = data[j];
  15.                         result[1] = j;
  16.                 }
  17.         }
  18. }
  19. // функция выполняемая на графическои процессоре
  20. __global__ void searchGPU ( float * data, int * result )
  21. {
  22.    int idx = blockIdx.x * blockDim.x + threadIdx.x;
  23.    //data [idx] = data [idx] + 1.0f;
  24.    if ( data [idx] == 2.0f )
  25.    {
  26.         result[0] = data [idx];
  27.         result[1] = idx;
  28.    }
  29. }
  30.  
  31. int main ( int argc, char *  argv [] )
  32. {
  33.   int n        = 16 * 1024 * 1024;
  34.   int numBytes = n * sizeof ( float );
  35.   int size_int = 2 * sizeof ( int );
  36.         //timer for host
  37.   clock_t start, finish;
  38.   double  duration;
  39.              // allocate host memory
  40.   float * a = new float [n];
  41.   int * host_res = new int[2];
  42.         //init host variables
  43.   for ( int i = 0; i < n-1; i++ )
  44.     a [i] = 0.0f;
  45.     a [n-1] = 2.0f;
  46.   host_res [0] = 0;
  47.   host_res [1] = 0;
  48.         //Выбираем способ расчета
  49.   printf("Select compute mode: 1 - CPU, 2 - GPU\n");
  50.   int mode;
  51.   scanf("%i", &mode);
  52.   if ( mode == CPU )
  53.   {     start = clock(); //start timer
  54.         searchCPU (a, host_res);
  55.         finish = clock(); //stop timer
  56.         duration = (double)(finish - start) / CLOCKS_PER_SEC; //time in seconds
  57.         printf( "time spent executing by the CPU: %.5f seconds\n", duration );
  58.   } else
  59.   if ( mode == GPU )
  60.   {      // allocate device memory
  61.     float * dev = NULL;
  62.         int * res = NULL;
  63.    
  64.     cudaMalloc ( (void**)&dev, numBytes );
  65.     cudaMalloc ( (void**)&res, size_int );
  66.  
  67.                    // set kernel launch configuration
  68.     dim3 threads = dim3(512, 1);
  69.     dim3 blocks  = dim3(n / threads.x, 1);
  70.                   // create cuda event handles
  71.     cudaEvent_t start, stop;
  72.     float gpuTime = 0.0f;
  73.     cudaEventCreate ( &start );
  74.     cudaEventCreate ( &stop );
  75.                  // asynchronously issue work to the GPU (all to stream 0)
  76.     cudaEventRecord ( start, 0 );
  77.         //copy data to device
  78.     cudaMemcpy      ( dev, a, numBytes, cudaMemcpyHostToDevice );
  79.     cudaMemcpy      ( res, host_res, size_int, cudaMemcpyHostToDevice );
  80.                //call device function
  81.     searchGPU<<<blocks, threads>>>(dev,res);
  82. //copy data to host
  83.     cudaMemcpy      ( host_res, res, size_int, cudaMemcpyDeviceToHost );       
  84.     cudaEventRecord ( stop, 0 ); // stop timer for device
  85.     cudaEventSynchronize ( stop );
  86.     cudaEventElapsedTime ( &gpuTime, start, stop ); //gpuTime - time in miliseconds
  87.                // print the gpu times
  88.     printf("time spent executing by the GPU: %.5f seconds\n", gpuTime/1000 );
  89.  // release resources (device)
  90.     cudaEventDestroy ( start );
  91.     cudaEventDestroy ( stop  );
  92.     cudaFree         ( dev   );
  93.   }
  94.   else
  95.   {  printf("not correct input");  return 1;   }
  96.      //print result
  97.      printf ( "value %i, index %i",host_res[0], host_res[1]);
  98. // release resources (host)
  99.     delete a;
  100.     delete host_res;
  101.     return 0;
  102. }

Время выполнения программы на ЦП составило 78 миллисекунд.
При выполнение на графическом процессоре, время выполнения программы составило 47 миллисекунд. Как видим прирост производительности порядка 40%. Не так уж много. Однако время измерения выполнения на графическом процессоре измерялось с учетом копирования исходных данных в память видеокарты и результата обратно в ОЗУ.
Результат выполнения программы на GPU без учета копирования данных составил около 9 миллисекунд. Таким образом, львиная доля времени используется на копирование данных в память GPU и обратно. Но, несмотря на это, все же достигается неплохой прирост производительности.

Программа выполнялась на платформе со следующими характеристиками:

Windows XP Professional (5.1, Build 2600) Service Pack 3
Processor: Intel(R) Pentium(R) 4 CPU 3.00GHz (2 CPUs)
Memory: 1024MB RAM

Display Devices
Name GeForce 8600 GT
Compute Capability 1.1
Clock Rate 1188 MHz
Multiprocessors 4
Warp Size 32
Regs Per Block 8192
Threads Per Block 512
Watchdog Enabled Yes
Threads Dimentions 512 x 512 x 64
Grid Dimentions 65535 x 65535 x 1

Memory InformationTotal Global 511.688 MB
Shared Per Block 16 KB
Pitch 256 KB
Total Constant 64 KB
Texture Alignment 256
GPU Overlap Yes

Performance InformationMemory Copy
Host Pinned to Device 3059.73 MB/s
Host Pageable to Device 1600.54 MB/s
Device to Host Pinned 3062.96 MB/s
Device to Host Pageable 1651.8 MB/s
Device to Device 5295.5 MB/s
GPU Core Performance
Single-precision Float 75641.8 Mflop/s
Double-precision Float Not Supported
32-bit Integer 15184 Miop/s
24-bit Integer 75647.5 Miop/s

Спасибо. Жду ответов и комментариев.

Forums: