программа поиска строки в файлах

ЭЭ ну вообщем хотел предложить желающим создать такое приложение
Думаю так как оно простое его можно будет сделать gpu-ОС независимым
Что полезно для так сказать развития и можно в дальнейшем использовать как шаблон для написания более сложных gpgpu программ

ТЗ
1.Должно работать из командной строки
2.Должно работать на картах АТИ (какие там начинают поддерживать stream ) и Нвидиа 8х ...
3.должно работать в unix - ax и в windows-ax
4.строка поиска ограничиваеться 256 байтами (это ограничение на параметры глобал функции в нвидии АТИ-шники могут этот параметр уточнить, т.к по мне и 64 байт с головой)

примеры использования

ggrep [stroka] [file_name_pattern] [dir] -[key]
ggrep [stroka] ищет относительно текущей директории во всех файлах включая поддериктории
ggrep [stroka] [file_name_pattern] ищет относительно текущей директории в файлах соответствующих pattern включая поддериктории

key= nsub - ключ который отменяет поиск в поддиректоритях

ну вообщем пишите ваши мнения

Forums: 

аргументирую

операция поиска подстроки в файле bandwidth bounded.
так что прирост производительности будет ниже прироста геморроя.

В то же время, если паттернов

В то же время, если паттернов много, то есть смысл (но оно уже будет CPU bound)

То ли в каком-то из GPU Gems, то ли я просто отдельно где-то читал, описывается антивирусный сканер с весьма приличным приростом производительности.

извините новичка

что значит bandwith bound если это по русски ограниченно полосой пропускания памяти(около 5 ГГ/с)
то что тут страшного цпу тоже ограничен полосой пропускания памяти (по моему что то около 6.4 Г/с)
но только проц не может делать одновременно множество операций сравнения
а жлу могет
для нвидии алгоритм думаю сделать такой
1)грузим в видеопамять все файлы в буфер (предположим 16Мб)
2)спускаем на этот буфер кернел который ест его с помощью операции исключ ИЛИ 16кб-порциями (размер шаред мемори)
3)оставшийся буфер уменьшаем с помощью какой-то хитрой редьюс функции

В ситуации, когда ищется одно

В ситуации, когда ищется одно регулярное выражение (пусть в длинном тексте и пусть нету очень далеких lookahead) вы будете ограничены скоростью, с которой этот текст запихивается в видеокарту. А это вовсе не 6-12 Гбайт/сек, как у современных CPU, а в разы меньше.

Если паттернов много, то конечно ситуация меняется.

а чего это меньше

видеопамять для того и создано чтоб пихать в нее мегабайты
даже спец высокоскоростной интерфейс АЖП для этого был придуман 8)

давайте посчитаем производительность такого поиска для для буфера 16М
(при этом будем считать что время загрузки в видеопамять и в оперативную память проца равны
и для каждой операции нужен один такт )

тогда для проца время выполнения поиска по массиву в 16М равна ~ 16М тактов

расчитаем такты для паралельного вычисления (ядро см ниже)
1)загрузка шаред мемори
16M/кол-во блоков/кол-во потоков * 2 такта
2)операция ХОR
16M/кол-во блоков/кол-во потоков*(длина_искомого_слова+2)
3)операция записи результата 600 тактов на каждый результат (вроде чэтого что то в доке от нвидии)
итого
16M/кол-во блоков/кол-во потоков * (2+2+длина искомого слова) + такты на запись результата

__constant__ char CONST[LEN] = "BLA-BLA-BLa";

struct LIST {
__global__ int offset,
__global__ LIST *next
} find;

void __global__ kernell( char *buf, char * CONST, int LEN,LIST find){
extern shared bank[];
//загружаем из видеопамяти в шаред мемори
int idx = BlockId.x * DimBlock.x + threadId.x;
bank[ threadId.x ] = buf[idx]
__syncthreadth();
//операция сравнения
char res = 0xFF;
for(int i=0;inext = new LIST(idx);
}
__syncthreadth();
}

Больше всего во всем этом меня смущает как записать результат
И не только время доступа к глобальной памяти, но и как сделать эту операцию так чтоб не произошло перезаписывание в одну ячейку глобальной памяти двух результатов

ядро коректно не отображалось сорри

__constant__ char CONST[LEN] = "BLA-BLA-BLa";
struct LIST {
__global__ int offset,
__global__ LIST *next
} find;

void __global__ kernell( char *buf, char * CONST, int LEN,LIST find){
extern shared bank[];
//загружаем из видеопамяти в шаред мемори
int idx = BlockId.x * DimBlock.x + threadId.x;
bank[ threadId.x ] = buf[idx]
__syncthreadth();

//операция сравнения
char res = 0xFF;
for(int i=0 i меньше LEN i++)
//сравним так например
res AND =bank[i] XOR CONST[i]

//запишем результат
if(res)
find.next = new LIST(idx)

syncthreadth()
}

Посчитаем выигрыш
16M
------------------------------------------------------------------------------------------------------------------------------------------
16M/кол-во блоков/кол-во потоков * (2+2+длина искомого слова) + такты на запись результата

= (если пренебречь тактами записи результата)

кол-во потоков * кол блоков
-------------------------------------------
4+длина искомого слова

Такты - это отлично все. Но в

Такты - это отлично все.

Но в реальной жизни, читая файл с диска, вы и будете ограничены скоростью диска, а не тактами.

P.S. для зарегистрированных пользователей есть возможность редактирования собственных комментариев, необязательно второй раз писать если не отобразилось корректно.

ну не всегда. скажем есть

ну не всегда.
скажем есть текстовый файл, и необходимо через него прогнать тысячу регулярных выражений.
тут скорость не упрется в чтение/запись файла.

Про ситуацию, когда паттернов

Про ситуацию, когда паттернов много - я выше уже написал. В этой ситуации перенос на GPU может иметь смысл.

угу про диск мы и забыли

программа поиска для цпу точно так же ограничена этой скоростью

я вижу трудность совсем в другом
найти то совпадения я найду, а как его записать
думаю надо городить какой-то массив результатов с переменной длиной в шаред мемори
если этот массив делать списком, то появляеться очень неприятная ситуация, когда 2 потока могут записать что то в одну ячейку

Результаты можно и в

Результаты можно и в глобальную память писать. Скажем, если их ожидается мало, то счетчик с атомарным инкрементом и, собственно, все.

Если много - ну значит побить выходную память на сегменты (по сегменту на тред), а атомарный счетчик - на выделение очередного сегмента.

сравните отношение операции

сравните отношение операции сравнения чисел (xor или ==) и чтения/записи в глобальную память (по тактам).
сосчитайте кол-во операций сравнения и операций чтения/записи.
делайте выводы.

я не говорю что будет медленнее, я говорю что цели не оправдывают средства.

сам и отвечу себе 8)

ф-ции FindFirstFile FindNextFile
Побегайло системное программирование под windows

вопрос еще оч интересный
программу надо писать так чтоб во время обработки ядром буфера н.1
паралельно в буфер н.2 грузилась инфа
после обработки буферы бы менялись местами

в документации вроде такая возможность(паралельная работа кернела и загрузка в видеопамять) описана, но для этого оперативная память проца должна быть page locked.
что ета за память такая я еще не разобрался
*)

режим dma для видяхи

ну вообщем нашел еще одну проблему
для реализации программы поиска надо сначала загрузить буфер в озу
а потом перегрузить этот буфер в видео озу

и тут у мя возник вопрос чисто теоритический к людям разбирающимся в архитектуре компов
а может для видях есть какой-то режим
когда видяха грузит данные прям с диска в видеопамять миную оперативу

а зачем? скорость чтения с

а зачем?
скорость чтения с винта несравнимо ниже скорости записи в ОЗУ видеокарты.

пс: такое невозможно.