Потоки.

[Aspexs]

Здравствуйте.
У меня есть некоторая проблема.
есть такой вот код(Аналог алгоритма Рабина - Карпа) :

int idx = blockDim.x*blockIdx.x + threadIdx.x;

if (idx <N-M)
{
for (int i = 0 ; i < M; i++)
{
t = (d*t + txt[idx + i])%q;
}
if ( p == t )
{
for (j = 0; j < M; j ++)
{
if (txt[idx+j] != pat[j])
break;
}
if (j == M)
{
cuPrintf("Pattern found at index %d \n", idx);
}
}
...............................................
}

Я обрабатываю файл в 20 мб.

Почему то я не могу запустить больше 2048 потоков, и это в сумме, не
важно на скольки блоках.
Видеокарта поддерживает 1024*1024*16 потоков и 65536*65536*65536
блоков.

И еще, код работает только в том случае, если вместо точек(в
коде),ставить какой либо cuРrintf.
А если не ставить, то от чего не пашет, не могу разобраться почему.

Заранее спасибо за помощь.

[Dmitry N. Mikushin]

Полный тест с ядром и хостовой функцией - в студию.

29 мая 2012 г., 2:40 пользователь Aspexs <alekse...@gmail.com> написал:

> --
> Вы получили это сообщение, поскольку подписаны на группу CUDA.CS.MSU.SU.
>
> Чтобы добавлять сообщения в эту группу, отправьте письмо по адресу cudac...@googlegroups.com.
> Чтобы отменить подписку на эту группу, отправьте сообщение по адресу cudacsmsusu...@googlegroups.com.
> О дополнительных функциях можно узнать в группе по адресу http://groups.google.com/group/cudacsmsusu?hl=ru.
>

[Vitaliy Chernov]

> Видеокарта поддерживает 1024*1024*16 потоков

Тут нет никакой ошибки? Судя по всему, речь идёт о видеокарте с CC 2.x, размерность блока (x, y, z) у неё действительно может быть до (1024, 1024, 64), но общее число тредов в блоке не больше 1024 (см. Table F-2 NVIDIA CUDA c Programming Guide).

29 мая 2012 г., 10:37 пользователь Dmitry N. Mikushin <maem...@gmail.com> написал:

[Aspexs]

код хост функции:
#include <stdio.h>
#include <time.h>
#include<string.h>
#include "cuPrintf.cu"
#define d 256
#define LEN 20* 1024 *1024
__host__ void Make_file()
{
FILE *fin = fopen("tmp.txt","w");
int i;
int k=0;

srand(time(0));

for (i = 0; i < LEN; i++)
{
if((!(rand() % 40000))) && (!(rand() % 1000)))
{
fprintf(fin,"49 49 2A");
k++;
printf("index %d \n", i);
}
else
{
fprintf(fin, "%c", 65 + rand() % 26);
}
}
printf("Pattern count %d \n", k);
fclose(fin);
}

код ядра:

char temp[2048];
int bl, N;
Make_file();
char *pat = "49 49 2A";
int M = strlen(pat);
FILE *text = fopen("tmp.txt","r+");
fpos_t pos;
int p = 0; // хеш значение для образца
int h = 1;
int i;
int q = 101;

for (i = 0; i < M-1; i++ )
{
h = (h*d)%q;
}

for (i = 0 ; i < M; i++)
{
p = (d*p + pat[i])%q;
}

char* devtemp;
char* devpat;

//Выделяем память для массива на видеокарте
cudaMalloc((void**)&devtemp, sizeof(char) * 2048);
cudaMalloc((void**)&devpat, sizeof(char) * strlen(pat));

//Хендл Евента
cudaEvent_t start,stop;
float gpuTime = 0.0f;

cudaEventCreate(&start); //Создаем
cudaEventCreate(&stop);

//cudaEventRecord(start, 0); //Записываем
//копируем данные в память видеокарты
cudaMemcpy(devpat, pat, sizeof(char) *
strlen(pat),cudaMemcpyHostToDevice);

cudaPrintfInit();
do
{
bl = fread(temp, 1, 2048, text);
N = strlen(temp);
cudaMemcpy(devtemp, temp, sizeof(char) *
2048,cudaMemcpyHostToDevice);
dim3 blockSize (1024,1); //размер используемого блока
dim3 gridSize(2,1); // размер используемого грида
cudaEventRecord(start, 0); //Записываем
searchGPU<<<gridSize,blockSize >>> (devpat, devtemp, N, M, q, p, h);
cudaEventRecord ( stop, 0 );
fgetpos(text, &pos);
pos -= 7;
fsetpos(text, &pos);
} while (bl == 2048);

cudaPrintfDisplay();
cudaPrintfEnd();

//cudaEventRecord ( stop, 0 );
cudaEventSynchronize ( stop );
cudaEventElapsedTime ( &gpuTime, start, stop );

printf("time spent executing by the GPU: %.5f seconds\n", gpuTime/
1000 );

cudaEventDestroy ( start );
cudaEventDestroy ( stop );
cudaFree ( devtemp );
cudaFree ( devpat );

fclose(text);
}
Как то так.

[Aspexs]

ДА, в блоке не больше 1024 потоков, но почему тогда я могу только 2
блока запустить?
или например при 512 потоках, только 4 блока.
Если больше, то программа не работает(

[Dmitry N. Mikushin]

Уважаемый Aspexs,

Полный код нужен не за тем, чтобы на него полюбоваться, а чтобы его
запустить и увидеть ошибки, на которые Вы указываете :) Поэтому прошу
Вас представить такой тест, который будет 1) сразу полностью
собираться 2) воспроизводить ошибку.

Скорее всего ошибка в самом ядре. Можно было бы сказать точнее, если
бы Вы привели код ошибки. В Вашем коде статус ошибок CUDA-вызовов
вообще не проверяется, это очень плохая практика работы и очень
хороший способ надолго зависнуть на решении непонятных проблем.

- Д.

29 мая 2012 г., 11:29 пользователь Aspexs <alekse...@gmail.com> написал:

[Aspexs]

#include <stdio.h>
#include <time.h>
#include<string.h>
#include "cuPrintf.cu"
#define d 256
#define LEN 20* 1024 *1024
__host__ void Make_file()
{
FILE *fin = fopen("tmp.txt","w");
int i;
int k=0;

srand(time(0));

for (i = 0; i < LEN; i++)
{
if((!(rand() % 40000))) && (!(rand() % 1000)))
{
fprintf(fin,"49 49 2A");
k++;
printf("index %d \n", i);
}
else
{
fprintf(fin, "%c", 65 + rand() % 26);
}
}
printf("Pattern count %d \n", k);
fclose(fin);

}

__global__ void searchGPU (char *pat, char *txt, int N, int M, int q,
int p, int h)
{
int j;
int t = 0;

int idx = blockDim.x*blockIdx.x + threadIdx.x;

if (idx <N-M)
{
for (int i = 0 ; i < M; i++)
{
t = (d*t + txt[idx + i])%q;
}
if ( p == t )
{
for (j = 0; j < M; j ++)
{
if (txt[idx+j] != pat[j])
break;
}
if (j == M)
{
cuPrintf("Pattern found at index %d \n",
idx);
}
}
}
}

int main()
{

вот такой код.
А на счет ошибок вы правы, надо бы их использовать.

[Dmitry N. Mikushin]

Не компилируется - несбалансированные скобки в if-условии и

aspexs.cu(117): error: no operator "-=" matches these operands
operand types are: fpos_t -= int

29 мая 2012 г., 15:10 пользователь Aspexs <alekse...@gmail.com> написал:

[Aspexs]

#include <stdio.h>
#include <time.h>
#include<string.h>
#include "cuPrintf.cu"


#define d 256

#define LEN 2039

__host__ void Make_file()
{
FILE *fin = fopen("tmp.txt","w");
int i;
int k=0;

srand(time(0));

for (i = 0; i < LEN; i++)
{

if (i == 15)//((!(rand() % 40000)) && (!(rand() % 1000)))

{
fprintf(fin,"49 49 2A");
k++;
printf("index %d \n", i);
}
else
{
fprintf(fin, "%c", 65 + rand() % 26);
}
}
printf("Pattern count %d \n", k);
fclose(fin);
}

__global__ void searchGPU (char *pat, char *txt, int N, int M, int q,
int p, int h)
{
int j;
int t = 0;
int idx = blockDim.x*blockIdx.x + threadIdx.x;

if (idx <N-M)
{
for (int i = 0 ; i < M; i++)
{
t = (d*t + txt[idx + i])%q;
}
if ( p == t )
{

/* Проверка символов по одному */

for (j = 0; j < M; j ++)
{

//cuPrintf("Pattern found at index %d \n", idx);

if (txt[idx+j] != pat[j])
break;
}
if (j == M)
{
cuPrintf("Pattern found at index %d \n", idx);
}
}
}

cuPrintf("\x00");

}
int main()
{
char temp[2048];

int bl;

Make_file();
char *pat = "49 49 2A";
int M = strlen(pat);
FILE *text = fopen("tmp.txt","r+");
fpos_t pos;
int p = 0; // хеш значение для образца
int h = 1;
int i;
int q = 101;

for (i = 0; i < M-1; i++ )
{
h = (h*d)%q;
}

for (i = 0 ; i < M; i++)
{
p = (d*p + pat[i])%q;
}

char* devtemp;
char* devpat;

//Выделяем память для массива на видеокарте

cudaMalloc((void**)&devpat, sizeof(char) * strlen(pat));


//Хендл Евента
cudaEvent_t start,stop;
float gpuTime = 0.0f;

cudaEventCreate(&start); //Создаем
cudaEventCreate(&stop);

//копируем данные в память видеокарты
cudaMemcpy(devpat, pat, sizeof(char) *
strlen(pat),cudaMemcpyHostToDevice);

cudaPrintfInit();
do
{
bl = fread(temp, 1, 2048, text);

int N = strlen(temp);
cudaMalloc((void**)&devtemp, sizeof(char) * N);
cudaMemcpy(devtemp, temp, sizeof(char) *
N,cudaMemcpyHostToDevice);

dim3 blockSize (1024,1); //размер используемого блока
dim3 gridSize(2,1); // размер используемого грида
cudaEventRecord(start, 0); //Записываем
searchGPU<<<gridSize,blockSize >>> (devpat, devtemp, N, M, q, p,
h);
cudaEventRecord ( stop, 0 );

cudaPrintfDisplay();
printf("time \n");
cudaFree( devtemp );
fgetpos(text, &pos);
pos = pos - 7;

fsetpos(text, &pos);

} while (bl == 2048);

cudaPrintfDisplay();
cudaPrintfEnd();

//cudaEventRecord ( stop, 0 );
cudaEventSynchronize ( stop );
cudaEventElapsedTime ( &gpuTime, start, stop );

printf("time spent executing by the GPU: %.5f seconds\n", gpuTime/
1000 );

cudaEventDestroy ( start );
cudaEventDestroy ( stop );

//cudaFree ( devtemp );

cudaFree ( devpat );

fclose(text);
}

Этот код Компилируется, но не выдает нужного результата.
А вот если брать по килобайту, и запускать 1024 потока, или 512
потоков и 2 блока, то работает.
От чего нельзя запустить больше 1024 потоков????