Подождите...Каталог
Новые концепции для CUDA. Часть 2Давайте посмотрим, как NVIDIA описывает среду программирования CUDA. Довольно наглядно это показано на следующем рисунке:
Термин 'kernel' (ядро ОС) взят от потоковых процессоров. CPU делегирует "ядро" на GPU, и оно там разбивается на сети, блоки и треды. Кроме того, на этом рисунке не показано, что треды (Thread) в блоках группируются в "четверки". На первый взгляд это может показаться нагромождением, но в действительности это не так. Четверки соответствуют группе тредов, которые работают синхронно, то есть в идеале в пределах одной четверки ветвление полностью когерентное, что обеспечивает максимальную производительность. Помните, что когерентная ветвь у G80 состоит из 32 тредов. Четверки также группируются в блоки. Идея здесь в том, что одиночный блок будет гарантированно исполняться на том же самом мультипроцессоре, что позволяет тредам в пределах одного блока поддерживать связь и синхронизироваться друг с другом. Программист может задавать количество тредов на блок – это влияет на допустимые задержки и максимальное число доступных регистров, а также на объем доступной общей памяти. Вернемся к этому чуть позже. После этого все блоки группируются в одну сеть, которая объединяет все треды, выделенные GPU под одно программное ядро. Программа GPU в любое время имеет доступ к уникальному идентификатору задействованного блока и треда. Если Вы хотите, чтобы одновременно работали несколько программных ядер (что действительно нужно в мульти-GPU конфигурации), нужно, чтобы CPU работал в многопоточном режиме. СинхронизацияСейчас вы, может быть, думаете, что для организации работы примитивов синхронизации потребуется уйма переменных и функций. Это не так – все это делается вызовом единственной функции, у которой даже нет параметров! __syncthreads() просто позволяет Вам устанавливать точку синхронизации – барьер, и до того, как все треды блока не закончат выполнение всех предыдущих инструкций, ни один последующий код не будет выполняться (блокирование кода, относящегося как к общей памяти, так и к памяти устройства). Это позволяет эффективно предотвратить всевозможные отказы общей памяти "чтение ещё до записи", а также памяти устройства, если другие блоки не работают с теми же самыми ячейками памяти. Теоретически, эта команда почти не заметна; она занимает не больше времени, чем любая другая одиночная скалярная операция (2 цикла для 32 тредов), но здесь есть небольшая загвоздка. Чтобы понять проблему, представьте себе, что только арифметико-логические устройства имеют около 10 ступеней (с точки зрения шедулера это 5 ступеней, так как он работает на половинной частоте). Когда Вы запускаете синхронизацию, Вы хотите сделать так, чтобы все треды были в одной и той же точке процедуры. Это означает, что вы сдвигаете конвейер! При сдвиге конвейера Вы теряете минимум столько рабочих тактов, сколько стадий в этом конвейере. И это даже без учета задержек, что ещё больше усложняет ситуацию. Если, например, считывание данных из памяти для одной из четверок не закончилось, то все другие четверки для поддержки синхронной работы будут ждать окончания этой операции! Понятно, что это плохо. Можно ли эту проблему полностью устранить? Нет. Можно ли устранить её хотя бы частично? Да. Основная идея в том, что если у вас для этой операции задействованы 2 блока на один мультипроцессор, можно надеяться, что у Вас ещё достаточно много других тредов, чтобы обеспечить приличное быстродействие, в то время как другие блоки засинхронизированы. Если, например, у вас работает 8 блоков на одном мультипроцессоре и синхронизации достаточно редки, то есть практически никогда не совпадают по времени, для того чтобы спрятать задержки памяти, у Вас все еще имеется 87,5% тредов. Так это работает для локальной синхронизации. Если Вам "без вариантов" нужна глобальная синхронизация, помните, что Вы можете просто распределить нагрузку на несколько программных ядер и заставить CPU осуществлять синхронизацию. Если в каком-то конкретном случае это сделать нельзя, или это будет неэффективно, то треды могут коммутировать друг с другом через память устройства, что очень похоже на то, как если бы задача решалась с помощью CPU. Только убедитесь в том, что Вы знаете, какие блоки в данный момент в вашей архитектуре GPU работают, так как число мультипроцессоров может быть разным! Методики предотвращения сбоев "чтение потом запись" в этом случае такие же, как и для CPU, хотя у последнего может быть задействован общий кэш L2. У CUDA вообще нет кэша с чтением/записью, поэтому для организации общей синхронизации нужно учитывать многочисленные задержки памяти. Очевидно, что в идеале нужно иметь быстродействующую локальную синхронизацию и глобальную синхронизацию, организованную на уровне кэша второго уровня, но, похоже, что с G80 мы этого ещё не имеем. [N5-Что в сумме] Теперь, когда мы знаем, что из себя представляет CUDA и что она может, давайте подведем краткие итоги:
Можно было бы и дальше развить её функциональными возможностями, например, ускорить не очень быстрый рендеринг, но можно возразить, что это бы напрасно усложнило API для тех, кому CUDA нужна не для рендеринга, а среди потенциальных клиентов таких людей большинство. Преимущества и целевые рынкиНа первый взгляд, CUDA и GPGPU, в общем, нацелены, прежде всего, на рынок науки. К сожалению, сейчас для этого ещё не все в порядке – вычисления с двойной точностью (FP64) не поддерживаются. Однако есть и хорошие новости – и NVIDIA, и AMD обещают сделать такую поддержку к концу этого года. Быстродействие при этом, конечно, понизится, но всё ещё будет очень хорошим. А для чего CUDA хороша уже сейчас? Оказывается, что для очень многого. Посмотрите на диаграмму, которую NVIDIA продемонстрировала на презентации G80:
Технические обоснования таких ускорений для разных областей применения очень разные. Для физики и волновых уравнений, весьма вероятно, такие данные определяются главным образом числом гигафлопов, которыми может похвастаться чип. Для биологических задач, если честно, объяснение дать затрудняемся. Матричные и финансовые расчеты тоже интересны. Можно бы было ёще посмотреть быстродействие библиотек FFT и BLAS у CUDA, но это выходит за рамки нашей статьи. Числовые матрицы очень хорошо работают с кэшем параллельных данных, а финансовые задачи выигрывают в скорости расчетов благодаря быстродействию специальных функций G80. Этому посвящена превосходная глава в книге GPU Gems 2 "Options Pricing on the GPU". Выигрыш в производительности при задействовании специальных функций очень большой, и на сегодня GPU являются и чрезвычайно быстрыми, и точными инструментами. Хотя, похоже, что здесь сравнивается быстродействие CPU и GPU при расчете с разной точностью, но и при более низкой точности расчетов на CPU отличие будет все равно ощутимым. Что касается численных матриц, то именно на примере умножения матриц NVIDIA показала эффективность кэша параллельных данных и синхронизации. Оказывается, что при умелом использовании этих факторов, можно значительно снизить требования к пропускной способности, что позволяет очень сильно увеличить производительность. Так как данные всех блоков матрицы загружаются в общую память одновременно всеми тредами, то ещё до начала вычислений нужно включить синхронизацию.
[N6-Имеющиеся ограничения и перспективы на будущее] Как уже говорилось выше, FP64 вычисления в настоящее время не поддерживаются ни одним доступным GPU, но в этом году такая поддержка должна быть реализована. Однако останется ряд других ограничений – вот список, в котором перечислены, вероятно, большинство из них:
Тем не менее, перспективы вычислений на GPU очень большие. В дополнение к дальнейшему усовершенствованию функций, нужно ожидать существенного увеличения производительности, которая в течение ближайших нескольких лет будет значительно опережать закон Мура. Значительное улучшение произойдет благодаря увеличению числа транзисторов и увеличения тактовой частоты, кроме того, увеличится процент транзисторов, участвующих в арифметических операциях. Хорошим примером здесь является R580, у которого, по сравнению с R520, число транзисторов и размер кристалла увеличился только на 22%, а число арифметико-логических устройств и размер регистрового файла увеличились в 3 раза. Скорее всего, и у G80 более высокое соотношение ALU, чем у R520, но, хочется верить, что в ближайшее время NVIDIA сможет это соотношение увеличить ещё в 2 или 3 раза. Если в архитектуре CPU не будет никаких революций, преимущество в производительности решений GPGPU, вероятно, в будущем будет только расти. ЗаключениеНедавно NVIDIA официально обнародовала CUDA, а также FFT и BLAS библиотеки для неё. Ещё рано делать выводы о программировании в ней и её быстродействии при работе с обычными алгоритмами, но такие обзоры, конечно же, будут. Надеемся, что этот обзор поможет Вам понять, что такое CUDA и каковы её возможности. Конечно, для тех, кто интересуется GPGPU, сейчас наступают волнующие времена, и, зная, что нас ждет в ближайшем будущем, можно только радоваться перспективам этой IT отрасли. Также очень интересно посмотреть, что здесь нам предложит AMD, когда наконец-то будет выпущен их процессор R600 с архитектурой следующего поколения. И наконец, пользуясь случаем, хотим попросить NVIDIA (и AMD), чтобы они не ограничивали искусственно развитие GPGPU на потребительских видеокартах. Понятно, что средние цены на рынке GPGPU будут расти, но хотелось бы напомнить, что для новых рынков нужны новаторские разработки. Если любому пользователю, имеющему некоторые навыки в написании программ и яркие идеи, дать возможность экспериментировать на реальном аппаратном обеспечении, это может стать катализатором инноваций в ближайшей и долгосрочной перспективе. Если сфокусироваться на дорогих продуктах, инноваций будет меньше и уменьшится привлекательность преимущества производительность/цена, которое делает текущие решения GPGPU такими интересными. Кроме того, такие ограничения полностью закроют возможность GPU-вычислениям стать полноправной частью связующих решений AAA игр и других неигровых потребительских приложений. Приложения, которые могли бы взаимодействовать с CUDA и CTM, а не только с DX10, должны иметь такую возможность. Хотя у нас нет никаких явных причин верить в то, что NVIDIA и AMD рассматривают вопрос об искусственном ограничении их GPGPU решений в потребительской сфере, это настолько важно с очень многих точек зрения, что об этом обязательно нужно сказать. Кроме уровня поддержки профессиональных решений, хотелось бы, чтобы, при дифференциации пользовательских продуктов от профессиональных, они ограничились только уровнем FP64, объемом памяти, функциональными возможностями мульти-GPU и другими подобными параметрами, а не набором функций. Большинство серверных CPU отличаются от других процессоров даже меньшим числом параметров, но от этого они только выиграли!
Источник: www.beyond3d.com | |||||||||||
