Алгоритм Рабина - Карпа

В общем тут такое дело, от чего то не могу запустить более 2048 нитей, и это в сумме на любом количестве блоков, в чем дело?
Пытаюсь нормально распараллелить алгоритм Рабина-Карпа.
Вот код, что с ним не так?
Если запускать например 1024 нити и 512 блоков, то он вообще функцию не выполняет(
Заранее спасибо за помощь.

  1. <br />
  2. #include <stdio.h><br />
  3. #include <time.h><br />
  4. #include<string.h><br />
  5. #include "cuPrintf.cu"<br />
  6. #define d 256<br />
  7. #define LEN 20 * 1024 * 1024<br />
  8. #define BLOCK_SIZE 512 * 1024<br />
  9. #define CPU 1<br />
  10. #define GPU 2</p>
  11. <p>__host__ void Make_file()<br />
  12. {<br />
  13.     FILE *fin = fopen("tmp.txt","w");<br />
  14.     int i;<br />
  15.     int k;<br />
  16.     srand(time(0));<br />
  17.     for (i = 0; i < LEN; i++)<br />
  18.     {<br />
  19.         if ((!(rand() % 109934)) && (!(rand() % 100)))<br />
  20.                 {<br />
  21.                         fprintf(fin,"49 49 2A");<br />
  22.                         k++;<br />
  23.                 }<br />
  24.                 else<br />
  25.                 {<br />
  26.                                          fprintf(fin, "%c", 65 + rand() % 26);<br />
  27.                 }<br />
  28.         }<br />
  29.         printf("Pattern count %d \n", k);<br />
  30.     fclose(fin);<br />
  31.     }<br />
  32. // функция выполняемая на CPU<br />
  33. __host__ void searchCPU (char *pat, char *txt,int N, int M)<br />
  34. {<br />
  35.     int  j;<br />
  36.     int q = 101;<br />
  37.     int p = 0; // хеш значение для образца<br />
  38.     int t = 0; // хеш значение для текста<br />
  39.     int h = 1;</p>
  40. <p>    for (int i = 0; i < M-1; i ++ )<br />
  41.     {<br />
  42.         h = (h*d)%q;<br />
  43.     }<br />
  44.     for (int i =  0 ; i < M; i ++)<br />
  45.     {<br />
  46.         p = (d*p + pat[i])%q;<br />
  47.         t = (d*t + txt[i])%q;<br />
  48.     }<br />
  49.     // Движение образца по тексту<br />
  50.     for (int i = 0; i <= N - M; i++)<br />
  51.     {<br />
  52.         // Проверяем хеш-значения текста и образца<br />
  53.         // Если хеш-значения совпадают, проверяем посимвольно<br />
  54.         if ( p == t )<br />
  55.         {<br />
  56.             /* Проверка символов по одному */<br />
  57.             for (j = 0; j < M; j++)<br />
  58.             {<br />
  59.                 if (txt[i+j] != pat[j])<br />
  60.                     break;<br />
  61.             }<br />
  62.             if (j == M)<br />
  63.             {<br />
  64.                 printf("Pattern found at index %d \n", i);<br />
  65.             }<br />
  66.         }<br />
  67.         // Вычисление следующего хеш-значения текста<br />
  68.         if ( i < N-M )<br />
  69.         {<br />
  70.             t = (d*(t - txt[i]*h) + txt[i+M])%q;<br />
  71.             // Мы можем получить отрицательное значение t, преобразцем в положительное<br />
  72.             if(t < 0)<br />
  73.             t = (t + q);<br />
  74.         }<br />
  75.     }<br />
  76. }<br />
  77. __global__ void searchGPU (char *pat, char *txt, int N, int M)<br />
  78. {<br />
  79.     int q = 101;<br />
  80.     int p = 0; // хеш значение для образца<br />
  81.     int t = 0; // хеш значение для текста<br />
  82.     int h = 1;</p>
  83. <p>    for (int i = 0; i < M-1; i ++ )<br />
  84.     {<br />
  85.         h = (h*d)%q;<br />
  86.     }<br />
  87.     for (int i =  0 ; i < M; i ++)<br />
  88.     {<br />
  89.         p = (d*p + pat[i])%q;<br />
  90.         t = (d*t + txt[i])%q;<br />
  91.     }<br />
  92.     // Движение образца по тексту </p>
  93. <p>     int idx = blockIdx.x * blockDim.x + threadIdx.x;<br />
  94.         // Проверяем хеш-значения текста и образца<br />
  95.         // Если хеш-значения совпадают, проверяем посимвольно<br />
  96.         if(p==t)<br />
  97.         {<br />
  98.                 int fMatch = 1;<br />
  99.                 /* Проверка символов по одному */<br />
  100.                 for (int j = 0; j < M; j ++)<br />
  101.                 {<br />
  102.                         if (txt[idx+j] != pat[j])<br />
  103.                                           fMatch = 0;<br />
  104.                 }<br />
  105.                 if (fMatch == 1)<br />
  106.                 {<br />
  107.                         cuPrintf("Pattern found at index %d \n", idx);<br />
  108.                 }<br />
  109.         }<br />
  110.         if ( idx < N-M )<br />
  111.               {<br />
  112.                       // Вычисление следующего хеш-значения текста<br />
  113.                    t = (d*(t - txt[idx]*h) + txt[idx+M])%q;<br />
  114.             // Мы можем получить отрицательное значение t, преобразцем в положительное<br />
  115.                    if(t < 0)<br />
  116.                    t = (t + q);<br />
  117.         }<br />
  118. }<br />
  119. int main()<br />
  120. {<br />
  121.     char temp[BLOCK_SIZE];<br />
  122.     int bl;<br />
  123.     Make_file();<br />
  124.     char *pat = "49 49 2A";<br />
  125.     int M = strlen(pat);<br />
  126.     FILE *text = fopen("tmp.txt","r+");<br />
  127.     fpos_t pos;<br />
  128.     int N;</p>
  129. <p>    //timer for host<br />
  130.     clock_t start, finish;<br />
  131.     double  duration;</p>
  132. <p>    printf("Select compute mode: 1 - CPU, 2 - GPU\n");<br />
  133.     int mode;<br />
  134.     scanf("%i", &mode);<br />
  135.     if ( mode == CPU )<br />
  136.     {<br />
  137.                 start = clock(); //start timer<br />
  138.                         do<br />
  139.                         {<br />
  140.                                 bl = fread(temp, 1, BLOCK_SIZE, text);<br />
  141.                                 N = strlen(temp);<br />
  142.                                 searchCPU (pat, temp, N, M);<br />
  143.                                 fgetpos(text, &pos);<br />
  144.                                 pos -= 7;<br />
  145.                                 fsetpos(text, &pos);<br />
  146.                         } while (bl == BLOCK_SIZE);</p>
  147. <p>             finish = clock();<br />
  148.                 duration = (double)(finish - start) / CLOCKS_PER_SEC; //time in seconds<br />
  149.                 printf( "time spent executing by the CPU: %.5f seconds\n", duration );<br />
  150.         }<br />
  151.         else<br />
  152.         if (mode == GPU)<br />
  153.         {<br />
  154.         //Указатели на память видеокарты<br />
  155.         char* devtemp;<br />
  156.         char* devpat;</p>
  157. <p>     //Выделяем память для массива на видеокарте<br />
  158.         cudaMalloc((void**)&devtemp, sizeof(char) * BLOCK_SIZE);<br />
  159.         cudaMalloc((void**)&devpat, sizeof(char) * strlen(pat));</p>
  160. <p>     //Хендл Евента<br />
  161.         cudaEvent_t start, stop;<br />
  162.         float gpuTime = 0.0f;</p>
  163. <p>     cudaEventCreate(&start); //Создаем<br />
  164.         cudaEventCreate(&stop);<br />
  165.         cudaEventRecord(start, 0); //Записываем</p>
  166. <p>     // initialize cuPrintf<br />
  167.   cudaPrintfInit();<br />
  168.         //Выполняем функцию ядра<br />
  169.         do<br />
  170.                 {<br />
  171.                         bl = fread(temp, 1, BLOCK_SIZE, text);<br />
  172.                         N = strlen(temp);<br />
  173.                                           //копируем данные в память видеокарты<br />
  174.                         cudaMemcpy(devtemp, temp,  sizeof(char) * BLOCK_SIZE,cudaMemcpyHostToDevice);<br />
  175.                         cudaMemcpy(devpat, pat,  sizeof(char) * strlen(pat),cudaMemcpyHostToDevice);</p>
  176. <p>                     dim3 blockSize (1024,1); //размер используемого блока<br />
  177.                         dim3 gridSize(2,1); // размер используемого грида</p>
  178. <p>                     searchGPU<<<gridSize,blockSize >>> (devpat, devtemp, N, M);</p>
  179. <p>                     fgetpos(text, &pos);<br />
  180.                         pos -= 7;<br />
  181.                         fsetpos(text, &pos);<br />
  182.                 } while (bl == BLOCK_SIZE);</p>
  183. <p>     // display the device's greeting<br />
  184.   cudaPrintfDisplay();<br />
  185.         // clean up after cuPrintf<br />
  186.   cudaPrintfEnd();</p>
  187. <p>     cudaEventRecord ( stop, 0 );<br />
  188.         cudaEventSynchronize ( stop );<br />
  189.         cudaEventElapsedTime ( &gpuTime, start, stop );</p>
  190. <p>     printf("time spent executing by the GPU: %.5f seconds\n", gpuTime/1000 );</p>
  191. <p>     cudaEventDestroy ( start );<br />
  192.     cudaEventDestroy ( stop  );<br />
  193.     cudaFree         ( devtemp );<br />
  194.     cudaFree         ( devpat );</p>
  195. <p>     }<br />
  196.         else<br />
  197.     {<br />
  198.     printf("not correct input");<br />
  199.     return 1;<br />
  200.     }<br />
  201.     fclose(text);<br />
  202. }<br />

Forums: