и кстати, в процессе решения одной задачи мы придумали забавный метод — наплевать на синхронизацию вовсе, а после вычисления свалидировать значения и где нужно подправить — конечно, специфика это.
не зная конкретной задачи могу только порассуждать о правильном использовании адресации, чтобы подзадачи пересекались по-минимуму.
есть такая интересная особенность, что потоки решающие одинаковую по размеру подзадачу обычно завершаются одновременно, этот факт можно косвенно использовать (например один идет с начала блока — второй с конца).
интересна примерная производительность i5/i7 в GFLOPS.
прогресс по обоим направлениям идет, однако пока видеокарты устойчиво выигрывают в отношении производительность/стоимость и производительность/энергопотребление. хотя, конечно, это может зависеть от задачи, в частности задачи, которые вообще не параллелятся, а одно ядро CPU все-таки значительно мощнее ядра GPU.
в общем-то спецификация __global__ позволяет выполнять одни и те же функции на CPU и GPU. но смысла в этом не очень много, мощностями CPU можно просто пренебречь.
а вот переписать код, который работает в рамках классической модели на видеокарты — задача, которая автоматически (пока) не разрешима — другая модель, все-таки.
интересно, а какое приблизительно получилось получить ускорение относительно CPU?
в задачах, которые мы решаем в основном очень большие объемы данных фигурируют, и единственная цель, написать их обработку, так чтобы время передачи на карточку было соизмеримо со временем работы алгоритма, так что конечно все эти 10-40 Гбайт/с больше, чем нужно. хотя многие алгоритмы работают в несколько проходов — там это важно.
Верно. Скажем так, за счет быстрой памяти можно эффективно обрабатывать очень большие объемы данных, при этом стараясь сделать так, чтобы обсчет не сильно испортил это время.
Про именно вычислительные задачи напишу потом, однако там не все так радужно.
в общем-то они и передаются (и используются, как же без них), но они не обязаны быть кратными размерности решаемой задачи, поэтому и передаю дополнительные параметры, которые можно как-то покрутить. в общем да, ничего серьезного, производительность отличается на единицы процентов.
+ некоторые if'ы сводятся к арифметическим выражениям. нет уверенности, что это именно так, но попытки их как-то оптимизировать вручную дали тот же результат что и исходная версия.
пробовал разные конфигурации, THREADS={96,160,256} x GRID={512,768,1024}, скорость примерно одинаковая, для других хуже, но это за счет не кратности данных задачи.
условия действительно очень простые, но ради развлечения составил if из 16 условий вида (x == 0 && y != 0 && z == 0) || (x != 0 && y == 0 && z == 0) ||…, они, конечно, без труда оптимизируются до 5-6 но на производительность как-то заметно не повлияло.
есть подозрение, что там что-то хитрее, чем просто параллельность по TRUE/FALSE.
По поводу GPGPU:
технология сильно совершенствуется, и сейчас нет больших проблем ускорить код который имеет большое количество ветвлений, а трансфер из памяти в видеопамять работает со скоростью прямого доступа к оперативной памяти (как более медленной). Вот как раз сегодня на тему написал: habrahabr.ru/blogs/hi/119435/
завтра попробую написать топик про то как мы это делали, и в частности проблема принадлежности точки к контуру была решена.
хотя, ради справедливости, в финальном релизе мы плюнули на оба варианта и реализовали «самый тупой» (алгоритм заливки) на CUDA. со скоростью обработки в 20Гбайт/с результирующее время даже не мерили.
и + скорость. мы когда занимались подобной задачей пользовали подобный (или в точности такой? уже не припомню) метод, время работы на картинках 3200x2400 измерялась в секундах. с внешними контурами — десятые доли секунды (реально сильно зависело от картинки, от 5 миллисекунд до секунды).
мой комментарий имеет смысл в предположении, что картинка в памяти доступна только для чтения.
собственно если это честный набор бит Ч/Б, то ее эффективнее представлять как набор бит, а если нет — то после работы алгоритма она не должна портиться.
контуры не надо убирать с изображения, достаточно просто списка пропусков точек для каждого столбца.
и да, если картинка поступает последовательно с построением контуров будут сложности.
данную задачу можно решить существенно быстрее и сильно сэкономив память.
сначала обходим объект по внешнему контуру (можно сделать за время порядка количества внешних точек, что обычно не больше 4*X*Y, если интересно, могу написать как делается), а потом заполняем соответствующим индексом все внутренние точки (если это нужно).
остается только уметь находить начальные точки для обхода, без предварительного квадратичного (X*Y) прохода по изображению это сделать сложно, но можно совместить этот процесс с считыванием изображения.
есть такая интересная особенность, что потоки решающие одинаковую по размеру подзадачу обычно завершаются одновременно, этот факт можно косвенно использовать (например один идет с начала блока — второй с конца).
прогресс по обоим направлениям идет, однако пока видеокарты устойчиво выигрывают в отношении производительность/стоимость и производительность/энергопотребление. хотя, конечно, это может зависеть от задачи, в частности задачи, которые вообще не параллелятся, а одно ядро CPU все-таки значительно мощнее ядра GPU.
несколько не понятно — а в чем проблема, если производится только чтение… или имелась в виду проблема с чтением/записью?
а вот переписать код, который работает в рамках классической модели на видеокарты — задача, которая автоматически (пока) не разрешима — другая модель, все-таки.
в задачах, которые мы решаем в основном очень большие объемы данных фигурируют, и единственная цель, написать их обработку, так чтобы время передачи на карточку было соизмеримо со временем работы алгоритма, так что конечно все эти 10-40 Гбайт/с больше, чем нужно. хотя многие алгоритмы работают в несколько проходов — там это важно.
Про именно вычислительные задачи напишу потом, однако там не все так радужно.
условия действительно очень простые, но ради развлечения составил if из 16 условий вида (x == 0 && y != 0 && z == 0) || (x != 0 && y == 0 && z == 0) ||…, они, конечно, без труда оптимизируются до 5-6 но на производительность как-то заметно не повлияло.
есть подозрение, что там что-то хитрее, чем просто параллельность по TRUE/FALSE.
технология сильно совершенствуется, и сейчас нет больших проблем ускорить код который имеет большое количество ветвлений, а трансфер из памяти в видеопамять работает со скоростью прямого доступа к оперативной памяти (как более медленной). Вот как раз сегодня на тему написал:
habrahabr.ru/blogs/hi/119435/
хотя, ради справедливости, в финальном релизе мы плюнули на оба варианта и реализовали «самый тупой» (алгоритм заливки) на CUDA. со скоростью обработки в 20Гбайт/с результирующее время даже не мерили.
собственно если это честный набор бит Ч/Б, то ее эффективнее представлять как набор бит, а если нет — то после работы алгоритма она не должна портиться.
контуры не надо убирать с изображения, достаточно просто списка пропусков точек для каждого столбца.
и да, если картинка поступает последовательно с построением контуров будут сложности.
сначала обходим объект по внешнему контуру (можно сделать за время порядка количества внешних точек, что обычно не больше 4*X*Y, если интересно, могу написать как делается), а потом заполняем соответствующим индексом все внутренние точки (если это нужно).
остается только уметь находить начальные точки для обхода, без предварительного квадратичного (X*Y) прохода по изображению это сделать сложно, но можно совместить этот процесс с считыванием изображения.
хотя вот майкрософтовская утилита fc (file compare) работает как-то по-другому и иногда дает очень странные результаты.