Свежие комментарии

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

Ты правильно описал путь к пизде###у - "варез программируют в сторонке три года и потом пытаются перевести на него продукты"

На самом деле переписывать с нуля надо так, чтобы цикл переписывания занимал максимум полгода (лучше - месяц) и все приложения тестировать/адаптировать в процессе переписывания. Если все за такой срок не переписывается - надо переписывать куски. Моя "мегаплатформе" такие "перепишем" уже переживала. Переписывал куски, в которых архитектура "сделаем по-быстрому, а потом начнем латать походу" начала сильно раздражать.

По поводу вчерашнего у Яндекса, Самохвалов 10000 потоков запустил?

Он там гонит (в другом месте) - криптуху на GPU программировали и до CUDA. В GPU Gems3 про AES пишут, но вообще довольно многое можно.

Другой вопрос, что с локальной памятью удобнее.

Вот Рамблер бы поднялся на эти бабки!

Я припоминаю, как у нас в Рамблере был администратор помешанный на rc5crack или как оно тогда называлось (не помню).

И я со своих серверов эту штуку выжигал каленой метлой....

Общая шина - это часто синдром третьей версии. Первая как-то работала, вторая - это куча патчей на первой, чтобы реально работать. В результате накопился опыт, денег достаточно (первые две версии - успешно продаются) и возникает желание сделать "по хорошему" - спроектировать все правильно, все предусмотреть и прочая.
Желание - вполне естественное, заметим.

Потом этот варез три года программируют, возможно в стороне от основного процесса, хуже когда отвлекая нужные ресурсы, а потом пытаются перевести на него продукты компании. Бывает часто, что и переводят с разрушительным эффектом. Бывает - просто палят бабки, это лучше и совсем не страшно.

Оно само проходит. В прошлые разы я ничего не делал для этого

Попробовал тупое копирование:
h_idata -- cudaMemcpy() --> d_idata
d_idata -- __global__ Copy() --> d_odata()
d_odata -- cudeMemcpy() --> h_odata
Результат у исхоника, приведёного ниже 50/50 (т.е. то работает, то не работает). Если уменьшать размер memSize - стабильность работы возрастает.

// includes, system
#include stdlib.h
#include stdio.h
#include string.h
#include math.h
#include time.h
#include windows.h

// includes, project
#include cutil.h
#include cuda.h

#define NUM_BLOCKS 1
#define NUM_THREADS 1

__global__ void Copy(char *from, char *to, int n)
{
for(int i=0; i "меньше" n; i++) to[i] = from[i];
to[0] = '1';
to[n-1] = '2';
}

int main(int argc, char** argv)
{
unsigned int memSize = 30000000;
FILE *F;

// host vars
char *h_idata = NULL;
char *h_odata = NULL;

// device vars
char *d_idata = NULL;
char *d_odata = NULL;

CUT_DEVICE_INIT(argc, argv);

//pinned memory mode - use special function to get OS-pinned memory
CUDA_SAFE_CALL( cudaMallocHost( (void**)&h_idata, memSize ) );
CUDA_SAFE_CALL( cudaMallocHost( (void**)&h_odata, memSize ) );

//allocate device memory
CUDA_SAFE_CALL(cudaMalloc((void**)&d_idata, memSize));
CUDA_SAFE_CALL(cudaMalloc((void**)&d_odata, memSize));

//initialize the memory
F = fopen("out.txt", "rb");
memSize = fread(h_idata, 1, memSize, F);
fclose(F);

//copy host memory to device memory

unsigned int timer = 0;
float elapsedTimeInMs = 0.0f;

CUT_SAFE_CALL( cutCreateTimer( &timer ) );
CUT_SAFE_CALL( cutStartTimer( timer));

// copy host to device
CUDA_SAFE_CALL(cudaMemcpy(d_idata, h_idata, memSize, cudaMemcpyHostToDevice));
// copy inside device
Copy "тут 3 меньше"NUM_BLOCKS, NUM_THREADS"тут 3 больше"(d_idata, d_odata, memSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
// copy device to host
CUDA_SAFE_CALL(cudaMemcpy(h_odata, d_odata, memSize, cudaMemcpyDeviceToHost));

CUT_SAFE_CALL( cutStopTimer( timer));
elapsedTimeInMs = cutGetTimerValue( timer);

printf("Elapsed: %f seconds\n", elapsedTimeInMs / (float)1000);

// write to file
F = fopen("output.txt", "wb");
fwrite(h_odata, 1, memSize, F);
fclose(F);

//clean up memory
CUDA_SAFE_CALL(cudaFreeHost(h_idata));
CUDA_SAFE_CALL(cudaFreeHost(h_odata));

CUDA_SAFE_CALL(cudaFree(d_idata));
CUDA_SAFE_CALL(cudaFree(d_odata));

CUT_EXIT(argc, argv);
}

В документации про глобальную память вроде ничего про ограничения не сказано... Что я не так делаю в этом простом копировании???

Т.е. у нас всетаки получается опросная модель? Или предлагается из select`а раз в секунду (цифра с потолка) вываливаться и пробовать захватить mutex (в каждом процессе)?

Мне кажется это плохая затея, в таком случае лучше использовать сценарий А, а еще лучше SCM_RIGHTS.

Алексей, вы общаетесь с Игорем, поэтому должны знать, у него воркеры на потоках или процессах?

В первом варианте меня как раз и смущают лишние wake up, все предыдущие посты на эту тему.

По второму... mutex отвечает за помещение хэндла серверного сокета в select?
Если так, то как раз возникает ситуация что остальные не узнают когда хэндл освободится если не будут постоянно молотить опрашивая mutex. А молотить мы не будем т.к. большую часть времени мы спим в select`е.

Все может быть гораздо хуже: у нормальной софтверной компании на начальном этапе все растет по экспоненте. Ну допустим, в два раза в год. Году на 4-5м может наступить полное видимое благолепие и запас устойчивости... на год-два. Или на 7-8-м году, видел такое.

Вот на фоне оного благолепия зарождаются монстры, которые за год-два могут съесть компанию изнутри. По очень большому числу причин, "общая шина" - это симптом, а не такой монстр. А причина - в том, что люди с экспоненциальной скоростью меняться обычно не могут, а при миллионном-многомиллионном обороте (и соответствующем размере компании со всеми партнерами, дилерами и прочая) - изменения быстро сделать тоже нельзя.

Алексей,
огромное спасибо за интересную статью!

У меня при работе с CUDA на 8800GTX возникли странные сложности с глобальной памятью :(

Задача: есть текстовый файл ~50Мб, есть словарь из 16384 слов, надо найти индексы первого вхождения каждого из этих слов, т.е. осуществить поиск. Естественно первое, что пришло на ум: запустить N блоков по M тредов, чтобы M*N=16384 параллельных тредов. Так и сделал. Но вот беда: при поиске в 50Мб выдаёт, что ни одного слова не нашёл (алгоритм - банальный брутфорс).

Тогда я сократил задачу и сделал 1 блок с 1 тредом в грайде (однопоточный поиск соответственно только 1-го слова из словаря, в файл в котором ищем вписал это слово в самый конец) - получилось то же самое :(
Затем я заменил в цикле брутфорса j=0 на j=40000000 и слово было найдено (производительность пока пофиг)!

Получается очень странно - по идее если мы выделим память с помощью cudaMalloc() и скопируем в неё из хоста - все потоки должны этот массив видеть _полностью_ , так ? А у меня все потоки видят только "по 10 мегабайт" (если слово поставить в пределах первых 10 Мб и вернуть j=0 - тоже находит).

Где могут быть "грабли" при работе с глобальной памятью?
Может есть какие ограничения?

P.S. Собсно вот исходник самого поиска для 1 треда:

//
// char *x - указатель на массив с искомым словом
// int m - длина массива x
// char *y - указатель на массив с текстом
// int n - длинна массива (50000000 байт в моём случае)
// int *r - указатель на массив с результатами (позиция в файле)
//
// все указанные выше массивы скопированы с хоста на девайс
// c помощью cudaMemcpy(), т.е. находятся в глобальной памяти девайса.
//
__device__ void BF(char *x, int m, char *y, int n, int* r)
{
int i, j;
*r = -1;

for (j = 0; j <= (n-m); j++)
{
for (i = 0; ((i < m) && ( *(x+i) == *(y+j+i) )); i++);
if (i >= m)
{
*r = j;
break;
}
}
}

//
// А это wrapper
//
__global__ void BFWrap(char *dict, int *offsets, char *y, int n, int* r)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int sid = bid * NUM_THREADS + tid;

if(offsets[sid+1]>0)
{
BF(dict+offsets[sid], offsets[sid+1] - offsets[sid] - 2, y, n, r+sid);
}
}

из main() запускаю его так:
BFWrap<<>>(d_ddata, d_odata, d_idata, textSize, d_rdata);

NUM_BLOCKS и NUM_THREADS за'define'ны и равны 1, т.е. я даже не пользую многпоточность и всё равно не работает :(