Комментарии 67
НЛО прилетело и опубликовало эту надпись здесь
Если топик вам понравится, то расскажу про то как…Топик нам понравился. Хочется действительно больше реальных применений, которые можно использовать в своих проектах. Пока что только кодирование аудио/видео припоминается.
Обработка изображений, фурье-анализ (хотя это туда же), решение систем уравнений, моделирование методом конечных элементов.
До учёных эта штука пока не очень докатилась, как докатится — начнём считать всевозможные коллапсы и прочее.
Я видел применения CUDA в научных работах с разных факультетов нашего университета. На факультете, где учусь я, основная проблема — отсутствие программистов. Я даже знаю лаборатории, где есть компьютеры с мощными карточками, но нету людей, которые с ними работают.
В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
> Я видел применения CUDA в научных работах с разных факультетов нашего университета.
Список публикаций, пожалуйста.
Даже больше скажу: scholar по запросу CUDA выдаёт только статьи с таким учёным, по запросу же OpenCL выдаёт статью-спецификацию OpenCL на которую уже есть 21 ссылка (за ~1 год)
> В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
И что-же они считают и как? Какие именно «циклотроны»?
Список публикаций, пожалуйста.
Даже больше скажу: scholar по запросу CUDA выдаёт только статьи с таким учёным, по запросу же OpenCL выдаёт статью-спецификацию OpenCL на которую уже есть 21 ссылка (за ~1 год)
> В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
И что-же они считают и как? Какие именно «циклотроны»?
Отвечу сам себе: FAIR использует CUDA: mpd.jinr.ru/trac/browser/trunk/cuda
Похоже используется для реконструкции трэков.
Похоже используется для реконструкции трэков.
Список я вам дать не могу, так как просто общался лично с людьми, это обсуждалось в неформальной беседе. Скажу лишь только то, что были это аспиранты и студенты, то есть речь шла о кандидатских и дипломы соответственно.
По Дубне можно найти тут информацию groups.google.com/group/cudacsmsusu/?hl=ru. Я с телефона, мне трудно найти ссылку сразу до нужного файла. Там где-то в прикрепленных файлах должна быть презентация.
По Дубне можно найти тут информацию groups.google.com/group/cudacsmsusu/?hl=ru. Я с телефона, мне трудно найти ссылку сразу до нужного файла. Там где-то в прикрепленных файлах должна быть презентация.
И ещё раз — нет, не только там:
www.iki.rssi.ru/seminar/20100225/perepelkin.pdf
Модель там, конечно, упрощённая. Но за наводку спасибо.
www.iki.rssi.ru/seminar/20100225/perepelkin.pdf
Модель там, конечно, упрощённая. Но за наводку спасибо.
Хотя вру — по CUDA есть две статьи, суммарное число ссылок порядка 500.
А как насчет рендеринга в 3DS Max? Можно ли CUDA приспособить под это и какой будет прок?
Спасибо. Сам полгода назад бодался с технологией на реальном коде и писал по итогам статью habrahabr.ru/blogs/hi/116579/
Были большие проблемы с ветвлениями, возможно последняя CUDA оптимизирована лучше. Впрочем, карта была послабее — gts450
А не мерили сколько насколько полученная производительность ниже теоретической?
И что профайлер на if ы показывает?
Были большие проблемы с ветвлениями, возможно последняя CUDA оптимизирована лучше. Впрочем, карта была послабее — gts450
А не мерили сколько насколько полученная производительность ниже теоретической?
И что профайлер на if ы показывает?
Профайлер все меряет на уровне функций. Если вы про число ветвлений, то их будет 0, поскольку столь маленькие if'ы считаются с помощью предикатов.
Ой, а можно насчет предикатов подробнее?
Есть несколько (по-моему 4) однобитных предикатных регистров. Соответственно команды сравнения могут записывать в них результат, а затем для любой другой команды можно указать, чтобы она исполнялась только если предикатный регистр установлен в false или true. Так что в любом случае исполняются обе ветки в if'e, но не создается нагрузки на шедулер (которому нужно в начале разбить варп, а потом как-то соединить).
+ некоторые if'ы сводятся к арифметическим выражениям. нет уверенности, что это именно так, но попытки их как-то оптимизировать вручную дали тот же результат что и исходная версия.
Спасибо, полезная статья.
И еще, насчет условий. Сколько kernelов запускалось параллельно?
Насколько я помню, любой if рубит производительность на 2 * время выполняемого под ним кода (сначала выполнятся одновременно kernelы на TRUE, потом на FALSE
Соответственно, ifы меньше влияют если
1) Под ними простая арифметика. Попробуйте поставить под условие что-то посерьезней одного действия из быстрой памяти
2) GPU недогружен, т.е. количество одновременно выполняемых kernelов меньше количества исполнимых блоков на карте.
Насколько я помню, любой if рубит производительность на 2 * время выполняемого под ним кода (сначала выполнятся одновременно kernelы на TRUE, потом на FALSE
Соответственно, ifы меньше влияют если
1) Под ними простая арифметика. Попробуйте поставить под условие что-то посерьезней одного действия из быстрой памяти
2) GPU недогружен, т.е. количество одновременно выполняемых kernelов меньше количества исполнимых блоков на карте.
пробовал разные конфигурации, 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.
условия действительно очень простые, но ради развлечения составил if из 16 условий вида (x == 0 && y != 0 && z == 0) || (x != 0 && y == 0 && z == 0) ||…, они, конечно, без труда оптимизируются до 5-6 но на производительность как-то заметно не повлияло.
есть подозрение, что там что-то хитрее, чем просто параллельность по TRUE/FALSE.
Однако, скорость в ваших примерах достигается в основном не за счет параллельности выполнения, а за счет быстрой памяти…
Иначе бы от количества одновременно использующихся вычислительных блоков зависела бы почти линейно.
Иначе бы от количества одновременно использующихся вычислительных блоков зависела бы почти линейно.
Верно. Скажем так, за счет быстрой памяти можно эффективно обрабатывать очень большие объемы данных, при этом стараясь сделать так, чтобы обсчет не сильно испортил это время.
Про именно вычислительные задачи напишу потом, однако там не все так радужно.
Про именно вычислительные задачи напишу потом, однако там не все так радужно.
Классическая хорошо ускоряемая на граф. ускорителе задача — молекулярная динамика.
Когда модель загружается в память ускорителя и там «трясется» сутками, довольно сложной математикой.
Я в основном с такими задачами встречался, и ваш пример очень интересен т.к. он в противоположен.
Когда модель загружается в память ускорителя и там «трясется» сутками, довольно сложной математикой.
Я в основном с такими задачами встречался, и ваш пример очень интересен т.к. он в противоположен.
интересно, а какое приблизительно получилось получить ускорение относительно CPU?
в задачах, которые мы решаем в основном очень большие объемы данных фигурируют, и единственная цель, написать их обработку, так чтобы время передачи на карточку было соизмеримо со временем работы алгоритма, так что конечно все эти 10-40 Гбайт/с больше, чем нужно. хотя многие алгоритмы работают в несколько проходов — там это важно.
в задачах, которые мы решаем в основном очень большие объемы данных фигурируют, и единственная цель, написать их обработку, так чтобы время передачи на карточку было соизмеримо со временем работы алгоритма, так что конечно все эти 10-40 Гбайт/с больше, чем нужно. хотя многие алгоритмы работают в несколько проходов — там это важно.
Ну, до самостоятельного написания молекулярной динамики я не дорос ещё. Но NVIDIA хвастается серьезным ускорением www.nvidia.com/object/tesla_bio_workbench.html
Через месяц будет Tesla, проверим на практике.
А в моей задаче (сложный ассоциативный поиск по большой базе данных) удалось сравниться с i5 на OpenMP. Но это, конечно, не тот результат, которого хотелось.
Там еще, кроме ветвлений, проблема конкурентного доступа к памяти на чтение.
Через месяц будет Tesla, проверим на практике.
А в моей задаче (сложный ассоциативный поиск по большой базе данных) удалось сравниться с i5 на OpenMP. Но это, конечно, не тот результат, которого хотелось.
Там еще, кроме ветвлений, проблема конкурентного доступа к памяти на чтение.
о, спасибо. как ни странно не сталкивался с готовой библиотекой (хотя, правда, она нам не очень поможет, но интересно).
несколько не понятно — а в чем проблема, если производится только чтение… или имелась в виду проблема с чтением/записью?
несколько не понятно — а в чем проблема, если производится только чтение… или имелась в виду проблема с чтением/записью?
Там блокировки по чтению одной ячейки несколькими потоками. Увы. В Fermi эта проблема менее остро стоит, но стоит. А в константную память не получается засунуть.
не зная конкретной задачи могу только порассуждать о правильном использовании адресации, чтобы подзадачи пересекались по-минимуму.
есть такая интересная особенность, что потоки решающие одинаковую по размеру подзадачу обычно завершаются одновременно, этот факт можно косвенно использовать (например один идет с начала блока — второй с конца).
есть такая интересная особенность, что потоки решающие одинаковую по размеру подзадачу обычно завершаются одновременно, этот факт можно косвенно использовать (например один идет с начала блока — второй с конца).
и кстати, в процессе решения одной задачи мы придумали забавный метод — наплевать на синхронизацию вовсе, а после вычисления свалидировать значения и где нужно подправить — конечно, специфика это.
И все-таки, если тест под рукой — попробуйте засунуть под пару ifов что-то поболее, например цикл.
Очень интересны результаты.
Очень интересны результаты.
>Количество потоков и размер сетки не влияет на итоговое количество вызовов в секунду
Ой-ой-ой.
Советую попробовать запустить пустой кернел в конфигурации GRID=1024, THREADS=512 и посчитать количество запусков в секунду. (будет сильно меньше)
Если надо, могу объяснить почему.
Ой-ой-ой.
Советую попробовать запустить пустой кернел в конфигурации GRID=1024, THREADS=512 и посчитать количество запусков в секунду. (будет сильно меньше)
Если надо, могу объяснить почему.
еще дополнение вместо передачи blockcount и blocksize в код кернела, можно использовать build-in переменные dim3 blockDim и dim3 gridDim
когда уже будут процессоры синхронизировать с видеокартами, т.е. потоки сами бы распределялись в зависимости от задачи/сложности )))
в общем-то спецификация __global__ позволяет выполнять одни и те же функции на CPU и GPU. но смысла в этом не очень много, мощностями CPU можно просто пренебречь.
а вот переписать код, который работает в рамках классической модели на видеокарты — задача, которая автоматически (пока) не разрешима — другая модель, все-таки.
а вот переписать код, который работает в рамках классической модели на видеокарты — задача, которая автоматически (пока) не разрешима — другая модель, все-таки.
Читал как-то в журнале Суперкомпьютеры, что есть проект поставить кластер под OpenCL.
Опять же, Intel уже выпустил альфу OpenCL для своих процессоров.
Было бы очень интересно.
А мощностями современных CPU пренебрегать не стоит, начиная с Nehalem они начали сокращать отставание с видеокартами.
Опять же, Intel уже выпустил альфу OpenCL для своих процессоров.
Было бы очень интересно.
А мощностями современных CPU пренебрегать не стоит, начиная с Nehalem они начали сокращать отставание с видеокартами.
интересна примерная производительность i5/i7 в GFLOPS.
прогресс по обоим направлениям идет, однако пока видеокарты устойчиво выигрывают в отношении производительность/стоимость и производительность/энергопотребление. хотя, конечно, это может зависеть от задачи, в частности задачи, которые вообще не параллелятся, а одно ядро CPU все-таки значительно мощнее ядра GPU.
прогресс по обоим направлениям идет, однако пока видеокарты устойчиво выигрывают в отношении производительность/стоимость и производительность/энергопотребление. хотя, конечно, это может зависеть от задачи, в частности задачи, которые вообще не параллелятся, а одно ядро CPU все-таки значительно мощнее ядра GPU.
Производительность CPU «отстаёт» от GPU. У CPU сейчас Тера-Флопы а у CPU «всего» Гига-Флопы. 5 Тера-Флопс (AMD 6990) vs 128 Гига-Флопс (Core i7). Проблем у GPU только две: PCI-Express и алгоритм. Не каждый алгоритм можно посадить на графическую карту. И одним из главных «бутылочных горлышек» это PCI-Express на графику и обратно в процессор. OpenCL хорош в том, что может использовать оба два типа устройств и хорош в том, что можно использовать «совместную» память как для CPU и GPU, но кернелы должны бы быть по идее разные для каждого типа.
Только не
Если нужны обе версии функции, то нужно писать сразу 2 спецификатора:
__global__. __global__ означает функцию, которая выполняется на GPU, а вызывается с CPU. Если нужны обе версии функции, то нужно писать сразу 2 спецификатора:
__device__ __host__. Компилятор сгенерирует две бинарные версии функции: одна будет вызываться и выполняться на CPU (этим правда займётся gcc/icc/cl), а другая делать тоже самое на GPU.Хотел бы немного уточнить то, как можно по-настоящему замедлить CUDA-программу.
То, что ты называешь «случайным» чтением, таким не является. Как раз-таки наоборот, в твоём примере самый быстрый вариант доступа – coalesced (когда все нити одного варпа обращаются к одному блоку памяти). Он в 8-16 раз быстрее по-настоящему случайного доступа. Вообще доступ к памяти (как на чтение, так и на запись) – самая медленная операция на GPU (200 и больше тактов). Другое дело, что если блоков много, то пока одни ждут – другие считают.
Нужно замедлить арифметические операции? Есть же остаток от деления! Вот где смерть производительности.
Логические операции? Сделай так, чтобы все нити варпа бежали по разным веткам условия ( if… else if… else if ...). Вот тогда просмотрим.
Справедливости ради хочется отметить, что CUDA CC 2.х устройства могут выделять память из пула с помощью malloc'ов…
Спасибо за статью! Почитать про структуры данных было бы интересно.
То, что ты называешь «случайным» чтением, таким не является. Как раз-таки наоборот, в твоём примере самый быстрый вариант доступа – coalesced (когда все нити одного варпа обращаются к одному блоку памяти). Он в 8-16 раз быстрее по-настоящему случайного доступа. Вообще доступ к памяти (как на чтение, так и на запись) – самая медленная операция на GPU (200 и больше тактов). Другое дело, что если блоков много, то пока одни ждут – другие считают.
Нужно замедлить арифметические операции? Есть же остаток от деления! Вот где смерть производительности.
Логические операции? Сделай так, чтобы все нити варпа бежали по разным веткам условия ( if… else if… else if ...). Вот тогда просмотрим.
Справедливости ради хочется отметить, что CUDA CC 2.х устройства могут выделять память из пула с помощью malloc'ов…
Спасибо за статью! Почитать про структуры данных было бы интересно.
Причём, я так понимаю, для чтения поток переводят на другой процессор.
А всевозможные if/else/for сейчас исполняются не очень оптимально, по крайней мере в ATI.
А всевозможные if/else/for сейчас исполняются не очень оптимально, по крайней мере в ATI.
Чтение происходит на другом процессоре, если оно идёт из текстур. Там хитрая схема преобразования адресов и данных. Если же программа обратилась к обычному адресу в RAM, то варп засыпает, происходит группировка обращений по соседним адресам и когда результат пришёл — скедулер снова переключается на этот варп.
Не знаю как в AMD, но в CUDA с ветвлениями всё двояко. Если правильно распределить данные по потокам, то ветвления не навредят. Циклы тоже работают неплохо, хотя если их разворачивать, то это даёт небольшой прирост производительности.
Не знаю как в AMD, но в CUDA с ветвлениями всё двояко. Если правильно распределить данные по потокам, то ветвления не навредят. Циклы тоже работают неплохо, хотя если их разворачивать, то это даёт небольшой прирост производительности.
Это вы по coalesced говорите, я правильно понимаю? Но ведь не все обращения такие?
Это общая схема. Если доступ к данным невозможно сгруппировать, то обращения полуварпа сериализуются и все эти потоки засыпают до тех пор, пока не приедут все результаты. То есть, если все 8 потоков полуварпа обратились к адресам, лежащим в 8 разных блоках памяти (1 блок – это выровненные по адресам 32/64/128 байтные сегменты в RAM), то делается 8 последовательных запросов к памяти.
Всё это хорошо, но где реальные программы использующие CUDA? Ну ладно, есть Премьер который использует ее. Всякие мутные шареваре-поделия типа МедиаЭспрессо и vReveal… А где массовые кодировщики видео/аудио/фото? Ниша практически пустая. Когда прикрутят к Megui, например?
CUDA – молодая технология. Далеко не все крупные компании готовы поставить на неё. Мало ли как завтра nVidia изменит API? Или вообще похоронит GPGPU на своих девайсах… Привязка к одному производителю видеокарт тоже играет роль.
Но некоторые таки внедряют. Matlab, Photoshop да и многие другие пакеты, где SIMD архитектура походит, используют видеокарты (хороший пример – перебор паролей).
Существуют реализации кодеков с использованием CUDA, но новые разработки с трудом опережают заоптимизированные насквозь реализации на CPU (с использованием SSE-инструкций и алгоритмических ухищрений).
Реальные примеры использования CUDA можно увидеть в научной среде. Я вот некоторые куски своих прог так и ускоряю.
Много примеров использования GPGPU можно увидеть в CUDA Zone
Но некоторые таки внедряют. Matlab, Photoshop да и многие другие пакеты, где SIMD архитектура походит, используют видеокарты (хороший пример – перебор паролей).
Существуют реализации кодеков с использованием CUDA, но новые разработки с трудом опережают заоптимизированные насквозь реализации на CPU (с использованием SSE-инструкций и алгоритмических ухищрений).
Реальные примеры использования CUDA можно увидеть в научной среде. Я вот некоторые куски своих прог так и ускоряю.
Много примеров использования GPGPU можно увидеть в CUDA Zone
Есть кроссплатформенный opencl (ATI, Cell, NVidia), поэтому к CUDA никто цеплятся не хочет.
В университете изучаем возможности применения CUDA для решения плохо обусловленных СЛАУ. Встала проблема: точность.
Двойная по стандарту IEEE754 поддерживается только с Compute Capability 1.3, и это пол-беды. Сам IEEE754 точность результата не описывает, а только точность операндов. Таким образом, берём плохо обусловленную СЛАУ, решаем в double extended числах, и получаем всего 17 знаков точности вместо «теоретических» 30+.
Все матпакеты валятся на наших тестах, и реализации BLAS на CUDA — не исключение. С производительностью — да, хорошо, а вот с точностью — хуже, чем на CPU.
Планируем писать библиотеку, строго обеспечивающую заданную точность элементарных операций. Пока в зачаточном состоянии, но, полагаю, можно будет «всего лишь» реализовать вычисления с произвольной точностью на CUDA.
Двойная по стандарту IEEE754 поддерживается только с Compute Capability 1.3, и это пол-беды. Сам IEEE754 точность результата не описывает, а только точность операндов. Таким образом, берём плохо обусловленную СЛАУ, решаем в double extended числах, и получаем всего 17 знаков точности вместо «теоретических» 30+.
Все матпакеты валятся на наших тестах, и реализации BLAS на CUDA — не исключение. С производительностью — да, хорошо, а вот с точностью — хуже, чем на CPU.
Планируем писать библиотеку, строго обеспечивающую заданную точность элементарных операций. Пока в зачаточном состоянии, но, полагаю, можно будет «всего лишь» реализовать вычисления с произвольной точностью на CUDA.
Попробуйте, для начала, переписать это на opencl.
Чем это будет лучше? Для nVidia, OpenCL – обёртка над CUDA. Причём IMHO более слабая в плане API, работающая медленнее и при этом менее удобная.
Таки плюс у неё в портабельности.
Работает медленней она процентов на 5-10 (читал где-то).
Стоит поменять карточку, посмотреть что будет на других картах.
Работает медленней она процентов на 5-10 (читал где-то).
Стоит поменять карточку, посмотреть что будет на других картах.
Выигрыш может составлять от 10 до 60 %… Как повезёт.
Писать на CUDA удобнее и быстрее – она более развесистая, ближе к плюсам (есть даже шаблоны и какой-никакой STL). OpenCL – гораздо более низкоуровневая (спасибо, AMD) и урезанная технология. Для меня скорость разработки и расчётов важнее «портабельности». Проще в требованиях к железу указать какая карточка нужна. Да, я не разрабатываю клиентского ПО и не занимаюсь GRID-вычислениями, так что могу себе позволить такую блажь)
Со временем, разрыв между технологиями уменьшится или вообще исчезнет. Для OpenCL появится и нормальный API и продвинутые библиотеки и хорошая литература. Но я как-то дожидаться этого не хочу. Лучше я уже сейчас буду пользоваться преимуществами нового железа.
Писать на CUDA удобнее и быстрее – она более развесистая, ближе к плюсам (есть даже шаблоны и какой-никакой STL). OpenCL – гораздо более низкоуровневая (спасибо, AMD) и урезанная технология. Для меня скорость разработки и расчётов важнее «портабельности». Проще в требованиях к железу указать какая карточка нужна. Да, я не разрабатываю клиентского ПО и не занимаюсь GRID-вычислениями, так что могу себе позволить такую блажь)
Со временем, разрыв между технологиями уменьшится или вообще исчезнет. Для OpenCL появится и нормальный API и продвинутые библиотеки и хорошая литература. Но я как-то дожидаться этого не хочу. Лучше я уже сейчас буду пользоваться преимуществами нового железа.
OpenCL — разработка Apple Inc. и Kronos Group. Или я чего-то не знаю?
А чем вы занимаетесь, если не секрет?
А чем вы занимаетесь, если не секрет?
В Kronos Group входят больше сотни компаний. Когда я говорил про вину AMD, я подразумевал, что у них очень низкоуровневый API для GPGPU. Когда я только начал заниматься вычислениями на видюхах у них вообще было что-то вроде распальцованного ассемблера (называлось CTM), сейчас у них есть более высокоуровневый инструмент — ATI Stream, кажется. Так как OpenCL — это кроссплатформенная прослойка над архитектурами (ATI Stream, nVidia CUDA, pthreads/win threads), то отталкивались они вероятно от наиболее жёсткого варианта. А больше всех привязан к особенностям железа именно AMD'шный вариант. Так что в итоге и получилось что у nVidia уже вовсю GPU приближается по возможностям к CPU (на Fermi есть общий кэш для всех потоков, можно динамически выделять память, запускать несколько кёрнелов на одном устройстве) + серьёзно обрастает библиотеками (хотя то, что cuRand появилась недавно и всем приходилось самим писать свои PRNG — этого я понять не могу) + обрастает API причём как низко-, так и высокоуровневым (CUDA 4.0 в этом отношении — достаточно серьёзный шаг вперёд, особенно для тех, кто разрабатывает MultiGPU приложения). А вот у OpenCL как-то всё кисло (хотя там я за обновлениями не особо слежу, но у них же по-прежнему API v 1.0 или 1.1?), за последние годы так никуда особо вроде и не продвинулись.
А занимаюсь я тем, что в аспере моделирую всякие процессы. В большинстве своём диффузию и структуры всякие. Для этого раньше писал софт практически полностью на GPU, а сейчас успокоился и уже более взвешенно подхожу к делу. Нахожу затыки в программах и переношу что возможно на CUDA. Занимаюсь этим уже почти 3 года, преподаю даже — рассказывааю студентам про С и Куду.
А занимаюсь я тем, что в аспере моделирую всякие процессы. В большинстве своём диффузию и структуры всякие. Для этого раньше писал софт практически полностью на GPU, а сейчас успокоился и уже более взвешенно подхожу к делу. Нахожу затыки в программах и переношу что возможно на CUDA. Занимаюсь этим уже почти 3 года, преподаю даже — рассказывааю студентам про С и Куду.
Мне почему-то кажется, что самый жёсткий вариант это не AMD даже, а pthreads/win threads.
Да, OpenCL v API 1.1.
Да, OpenCL v API 1.1.
Ну значит за последние полтора года ничего сильно не поменялось).
Дело в том, что CPU-потоки не накладывают практически никаких ограничений на распараллеливаемые куски кода. А вот GPU жёстко закрепляют список допустимых типов аргументов функций, делает всякие деления памяти на shared/global и прочие интересные вещи типа ограничений на размерности блоков/гридов. nVidia тщательно старается это всё скрыть и сгладить (иногда удачно, иногда не очень). Сейчас даже ООП можно применять на CUDA (если это кому-то надо)), серьёзно проработаны самые тормозные моменты, чтобы типичные для CPU операции не тормозили так сильно на GPU.
Кроме того CPU-потоки это полновесные процессы, которые делают сложные операции и практически не влияют друг на друга (ибо MIMD). Так как GPU – это SIMD, то нужно запускать потоки даже не по количеству процессоров, а с запасом в пару порядков и близко бегущие потоки будут сильно взаимодействовать. Поэтому я вообще слабо понимаю мотивацию для затраты кучи сил и времени на разработку такой технологии, с помощью которой можно работать на CPU также неудобно как на GPU). Лучше бы допилили GPGPU технологии.
Дело в том, что CPU-потоки не накладывают практически никаких ограничений на распараллеливаемые куски кода. А вот GPU жёстко закрепляют список допустимых типов аргументов функций, делает всякие деления памяти на shared/global и прочие интересные вещи типа ограничений на размерности блоков/гридов. nVidia тщательно старается это всё скрыть и сгладить (иногда удачно, иногда не очень). Сейчас даже ООП можно применять на CUDA (если это кому-то надо)), серьёзно проработаны самые тормозные моменты, чтобы типичные для CPU операции не тормозили так сильно на GPU.
Кроме того CPU-потоки это полновесные процессы, которые делают сложные операции и практически не влияют друг на друга (ибо MIMD). Так как GPU – это SIMD, то нужно запускать потоки даже не по количеству процессоров, а с запасом в пару порядков и близко бегущие потоки будут сильно взаимодействовать. Поэтому я вообще слабо понимаю мотивацию для затраты кучи сил и времени на разработку такой технологии, с помощью которой можно работать на CPU также неудобно как на GPU). Лучше бы допилили GPGPU технологии.
К сожалению, в графических задачах точность не особо важна, поэтому с ней на GPU пока туго.
Улучшить это дело – задача интересная и почётная. Только думаю придётся очень серьёзно покопаться в PTX, чтобы писать действительно быстрый код с большим количеством побитовых операций.
Улучшить это дело – задача интересная и почётная. Только думаю придётся очень серьёзно покопаться в PTX, чтобы писать действительно быстрый код с большим количеством побитовых операций.
Код reduction (min/max) похож на код с сайта AMD, переписанный для CUDA ;-)
Книжка сильно в тему:
www.twirpx.com/file/413014/
www.twirpx.com/file/413014/
Зарегистрируйтесь на Хабре, чтобы оставить комментарий
CUDA: аспекты производительности при решении типичных задач