GPGPU.ru

Административное: call for moderator

lexa — Sat, 26 Jun 2010 15:53:04 +0000

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

Летом я не всегда могу его оперативно вычищать, ищу помощника(ов), которые могли бы раз в 1-2 дня просматривать все поступившие свежие комментарии и на спам жать кнопку "это спам".

Предпочтение отдается давно зарегистрированным постоянным писателям на сайт.

Пишите в обратную связь или просто по почте (lexa at lexa dot ru) с указанием своего аккаунта на gpgpu.ru

OpenCL 1.1

lexa — Mon, 14 Jun 2010 16:01:07 +0000

Новостные сайты массово написали, что Khronos Group выпустила спецификации OpenCL 1.1 (вот, скажем, HardOCP).

Удивительно, но на сайте Khronos об этом знаменательном событи не написано (хотя вот на HardOCP лежит 3-страничный пресс-релиз). Ну да новость сегодняшняя, сначала пресс-релиз, потом сайт...уже написано, за час справились. К сожалению, удобного содержательного диффа не обнаружилось (кроме пресс-релиза), чтобы понять разницу - нужно 300+ страниц спецификации честно прочитать.

Из пресс-релиза, естественно, ничего существенного понять нельзя.

Но есть в этом и содержательная сторона: NVidia моментально прислала письмо, что у них есть Conformance Candidate (в письме, впрочем, в некоторых местах написано что OpenCL 1.0, ну так торопились), который доступен на девелоперском сайте.

И действительно, доступен. В числе прочего, там лежат виндовые драйвера версии 258.19, которые, как я понимаю, новее и текущих доступных на оф-сайте бета-драйверов 257.15 и прилагавшихся к свежему Нексусу 257.19.

А значит можно потестироваться на предмет, исправились ли странности бета-драйверов 257.15.

CUDA 3.1 beta

lexa — Sun, 16 May 2010 05:39:45 +0000

На девелоперском сайте NVIdia (увы, нужен логин туда), выложили CUDA 3.1 Beta.

Из реально вкусного, там 16-way kernel concurrency, плюс обновления для CUBLAS/CUFFT.

Но:

Примеров (GPU Computing SDK) пока нету, только сам CUDA Toolkit.
Драйвера, поддерживающие все это дело - только для Tesla. Наверное, можно их поставить обычным хаком (прописав PCI IDs в setup.inf), но я не стал. Update: драйвера 257.15 для GTX480 появились на сайте NVidia в разделе с бета-драйверами.

Релиз обещают "в следующем месяце", проще подождать.

Update: в форум запостили ссылку на TPB со всем этим счастьем. Update2: С драйверами 197.xx не работает.

Алгоритмы поиска на графическом процессоре с помощью технологии NVidia CUDA

lex_NTU — Thu, 29 Apr 2010 19:10:38 +0000

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

Введение

Данная статья является продолжением темы «CUDA search». В этой теме рассматривался простейший линейный поиск на графическом процессоре с помощью технологии nVidia CUDA. В наихудшем случае т.е. когда искомый элемент массива был последним, алгоритм на GPU превзошел свой аналог на CPU почти на 40%. Количество элементов в массиве было равно количеству процессоров. Такие условия задачи достаточно хороши для GPU. Однако, как быть в других условиях? Как с ними справиться GPU? Подходит ли он для других алгоритмов? Что ж, попытаемся ответить на эти вопросы.

Бинарный поиск для CPU

Почему бинарный? Во-первых: классическая его реализация достаточная проста в программировании. Другие модификации и улучшения могут оказаться полезными лишь при определенных условиях. Мы же рассматриваем общий случай. Во-вторых: «максимум за величину порядка ln2 N сравнений мы найдем искомый ключ (либо установим, что его нет в таблице)» [1, 440c]. Не много других алгоритмов смогут потягаться с таким быстродействием. И так имеем в самом худшем случае ln2 N сравнений. Правда алгоритм применим лишь для отсортированного списка. Реализация на CPU имеет следующий вид:

Исходный код функции выполняемой на CPU:

__host__ void searchCPU ( float * data, int * result )
{
   int           x =  16777200; // искомый элемент
   int i = 0; // левая граница массива
   int j = 16 * 1024 * 1024; // правая граница    массива
   int k = 0; 
   while ( i <= j )        
   {
      k = i +  (j - i) / 2 ; 
      if ( data[k] <  x )
         { i = k + 1; } 
      else
         if ( data[k] >  x )
            { j = k - 1; }
      else
      {
         result[0] = data [k];
         result[1] = k;
         break;
      }              
   }              
} 

Как видим, действительно ничего сложного. Данный алгоритм значительно превосходит по быстродействию простой линейный поиск (во много сотен раз), реализованный даже на GPU, ведь для линейного поиска имеем N сравнений в худшем случае. Но, как уже говорилось выше, он возможен лишь для отсортированного списка (о сортировке поговорим чуть ниже). Кстати говоря, этот алгоритм намного устойчивее, так как для него практически не имеет значение в какой части списка находиться элемент, тогда как для линейного - от этого напрямую зависит время поиска. Но это в реализации на CPU! На GPU для линейного поиска имеем противоположный случай т.е. достаточно устойчивый алгоритм, т.к. множество потоков может почти одновременно охватить различные участки списка.

Бинарный поиск для GPU.

На попытку написать бинарный поиск меня вдохновила статья о параллельном поиске на видеокартах, в которой приводиться график сравнения производительности для CPU и GPU [2], который изображен на рисунке 1.

Рисунок 1 - Absolute performance of index search implementations exploiting thread level and SIMD parallelism on the CPU and the GPU with respect to workload.

Как видим из графика, быстродействие GPU начинает превосходить свой аналог на CPU лишь на определенной отметке, т.е. при достаточны больших объемах дынных. Попытаемся реализовать бинарный поиск для GPU и подтвердить или опровергнуть данные в вышеуказанном графике. Но каким образом разделить вычисления между потоками? Первая мысль была самая простая. Разделить весь список на равные участки, в каждом участке выполнять бинарный поиск. Поспешно написанная программа дала невероятный результат. А именно, очень медленный. Нужно заметить, что чем меньше потоков запускалось, тем быстрее выполнялась программа, т.е. ее быстродействие стремиться к одному потоку. Размышления подтвердили бессмысленность написанного. Подумайте сами, например, имеем список из 100 000 элементов и 1 000 потоков. Разделяем список на потоки, получается, что для каждого потока имеем по 100 элементов. Действительно для 100 элементов бинарный поиск должен выполниться быстрее, чем для 100 000. Но как же другие 999 потоков? Они изначально обречены на неудачу. Они будут безуспешно пытаться найти искомый элемент. А сколько ресурсов уйдет на их запуск и выполнение? Нужно искать другое решение. После теоретических размышлений оно найдено! Нужно реализовать как бы не бинарный поиск, а поиск такого порядка, сколько у нас потоков. Например, имея 32 потока, будет 32-нарный поиск.
Алгоритм таков: делим весь список на количество потоков, Каждый поток смотрит один элемент и сравнивает его с искомым. Если это искомый элемент замечательно. Если текущий элемент меньше, запоминаем его как нижнюю границу. Если больше запоминаем как верхнюю границу. Каждый из потоков найдет свой элемент больший или меньший искомого, на нам нужны только граничные элементы, между которыми лежит искомый, следовательно, нужно среди меньших элементов оставить максимум, а среди больших – минимум. Далее повторяем процедуру, но уже не для всего списка а для новых его границ. Повторяем до тех пор пока либо не найдем элемент, либо количество элементов между нижней и верхней границей станет меньше количества потоков. В таком случае для оставшихся элементов применяем любой другой алгоритм, например однопоточный бинарный поиск. Рассмотрим пример.
Пусть имеем 2048 элементов и 32 потока. Для удобства индексы элементов в массиве и их значения будем считать одинаковыми. Ищем, скажем, элемент под номером 65 равный 65. И так первый шаг. Для каждого потока смотрим каждый 32 элемент. т.е. первый поток смотрит первый элемент, второй – элемент номер 32, третий элемент номер 64 и т.д. Каждый поток находит либо больший либо меньший элемент. В данном случае меньшие элементы 1, 32, 64. большие элементы 96, 128...2016. Среди меньших элементов максимальным является 64, среди больших минимальным является 96. И так имеем новые границы списка 64..96. Количество элементов между границами равно 96-64 = 32. Повторяем итерацию. 32 элемента делим на количество потоков, каждый поток смотрит свой элемент. Получает что каждый поток смотрит по одному элементу и находит искомый. Итого имеем всего два шага! Теоретически неплохо. Действительно если оценить алгоритм теоретически, то имеем скорость O = (ln количество потоков N + ln2 количество процессоров) в худшем случае. Теоретически очень неплохо. Попробуем реализовать.

Исходный код возможной реализации, функция для GPU:

__global__ void searchGPU ( float * data, /*float * search,*/ int * result )
{
   int x = 16 * 1024 * 1024; // количество элементов
   int mx = 0; // нижняя граница
   int mn = x-1; // верхняя граница 
   int step; // шаг
   x =  16777200; // искомый элемент
   int idx; // текущий поток
   while ( mn - mx > 32 ) // пока количество элементов между границами 
      //больше количества потоков (в данном случае потоков 32)
    {
      step = (mn - mx) / 32;                                          
      nt idx = mx + (blockIdx.x * blockDim.x + threadIdx.x) * step;                     
      if ( data[idx] == x ) // найден искомый элемент
         {
            result[0] = data [idx];
            result[1] = idx;
            break;  
         }else
         if ( data[idx] > x )
            {
               if(data[idx] < mn)
               mn =  data[idx]; // минимальный среди больших
         }else
         if(data[idx] < x)
         {
            if(data[idx] > mx)
            mx =  data[idx]; // максимальный среди меньших
         }                                                                            
      }
} 

Теоретический анализ говорит, что алгоритм должен иметь очень хорошую скорость, в действительности имеем скорость почти в 100 раз большую, по сравнению с первой реализацией, однако, увы, все еще намного меньшую скорости бинарного поиска для CPU.
В чем же дело? Причин много. Например, копирование данных на видеокарту и обратно. Ресурсы на запуск потоков. Задержка при произвольном чтении элементов из памяти. Все эти факторы сложно учитывать при теоретическом анализе алгоритма. Однако вполне возможно, что реализация данного алгоритма превзойдет бинарный поиск для CPU в задачах внешнего поиска. Когда поступает очень много данных с внешнего накопителя, а не с оперативной памяти. Но тут нужно учитывать много других факторов, например скорость чтения данных с устройства, организацию хранения самих данных на устройстве (да и данные еще должны быть отсортированы), и ориентировать алгоритм уже под эти условия.
Итак, алгоритм бинарного поиска не очень хорош в реализации на GPU. Но это в случае, если мы ищем один элемент в массиве. А если нам нужно найти массив элементов в массиве? Такая задача идеально подходит для GPU. Тогда мы организуем алгоритм таким образом, что бы каждый поток искал свой элемент, что будет происходить почти одновременно. Теоретически имеем все туже скорость - O = ( ln2N ). Тогда как для CPU реализация представляет из себя последовательный двоичный поиск каждого элемента, что составит O = ( ln2N * M ), где M – количество искомых элементов. Попробуем проверить сказанное на практике

Фрагмент программы, функция для GPU:

__global__ void searchGPU ( float * data, float * search)
{ 
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   //data – массив данных
   //search – искомый массив
   int i = 0;
   int j = 16 * 1024 * 1024;                      
   int k = 0;
   while ( i <= j )        
   {
      k = i +  (j - i) / 2 ;
      if ( data[k] <  search[idx] )                 
          {
              i = k + 1;                
         } 
      else
         if ( data[k] >  search[idx] )
            {              
               j = k - 1;  
            }
      else
           {
              break;     //элемент найден               
           }                             
   } 
} 

Данная программа на GPU работает в тысячи раз быстрее, чем на CPU, что подтверждает вышесказанное об эффективности алгоритма на GPU в случае поиска массива элементов.

Сортировка на GPU.

Какую из сортировок лучше всего выбрать для реализации на GPU? «Было изобретено достаточно большое количество алгоритмов сортировки, основных из них более 25 «такое пугающее количество методов на самом деле лишь малая толика всех алгоритмов, придуманных на сегодняшний день. Почему же существует так много методов сортировки? Ответ если не очевиден, то в полнее понятен – каждый метод имеет свои преимущества и недостатки, поэтому он оказывается эффективнее других при некоторых конфигурациях данных и аппаратуры» [1, 76c].
Обратимся за помощью к книге Кнута, «искусство программирования» в которой он проводит достаточно подробный анализ всех алгоритмов и подводит резюме.
«Алгоритм битонной сортировки, полезен, если можно одновременно выполнять большое число сравнений» [1, 400c] Т.е. алгоритм идеально подходит для реализации на GPU. А теперь Взглянем в SDK и находим там готовый пример, реализации данного алгоритма на CUDA.
Битонная сортировка, несомненно, является быстрее её аналога для CPU. Таким образом, я считаю, что если в задаче поиска имеется подзадача сортировки, то вполне оправданно использовать сначала битонную сортировку на GPU (хотя это так же зависит от конфигурации данных и аппаратуры), а затем применить бинарный поиск на CPU (так же зависит конфигурации данных и аппаратуры). Однако, это в тех случаях, когда сортировка окупит себя в дальнейшем. Если постоянно ведется поиск данных, то разумеется, имеет смысл упорядочить эти данные, что бы потом сэкономить время на поиске.

Снова о поиске

Возможны и другие случаи, например, когда данные динамически меняются, или просто поиск осуществляется относительно редко. Отличным примером может служить обычный текст. Возможность поиска в текстовом редакторе неотъемлемая его функция. В данном случае, имеет смысл организовать поиск с помощью параллельной технологии т.к. однопоточный алгоритм на CPU вынужден будет просматривать весь текст последовательно (пусть и не по каждой букве, а с определенным шагом), что к тому же, делает алгоритм неустойчивым, т.к. время поиска напрямую зависит от положения искомой строки в тексте. В случае же алгоритма на GPU имеем более независимый поиск, т.к. множество потоков может начинать просмотр текста с различных его участков.
Не будем подробно рассматривать этот случай, так как он будет аналогичен поиску элемента в числовом массиве, рассмотренному в первой статье. В наилучшем и среднем случае алгоритм на CPU безусловно выиграет по времени поиска, однако в худшем случае, когда искомый элемент будет в конце списка, он проиграет GPU. Разумеется, если речь идет о достаточном большом количестве данных. Какой из алгоритмов лучше использовать однозначно ответить нельзя. В случае CPU мы будем иметь поиск, который может значительно растягиваться по времени, а в случае с GPU будем иметь достаточно постоянное время поиска.

Выводы

И так, подведем итоги. В общем можно сказать, что GPU особенно приспособлены решать проблемы, которые могут быть адаптированы к параллельным вычислениям, когда одна и та же программа (алгоритм) выполняется над многими элементами данных параллельно, требующих интенсивных расчетов: от простых арифметических операций до операций над памятью. Выполнение одной и той же программа для многих элементов данных одновременно снижает требование к изощренному контролю потоков.
Параллельное вычисление распределяет элементы данных между одновременно выполняющимися потоками. Многие приложения, которые обрабатывают большие объемы данных могут использовать модель параллельных вычислений для повышения производительности. Многие алгоритмы не из области обработки изображений могут быть улучшены с помощью параллельных вычислений, начиная с обработки сигналов или физического моделирования систем и заканчивая вычислительной экономикой или биологией.
Однако ускорение алгоритмов с помощью технологии nVidia CUDA далеко не панацея и имеет ряд недостатков и ограничений среди которых особенности архитектуры и особенности программирования. Теоретический анализ алгоритмов не всегда применим в случае технологии nVidia CUDA, т.к. не учитывает многих ее особенностей. Например, вложенный механизм распараллеливания функций и запуска потоков, особенности работы с памятью, конфигурацию оборудования и .т.п. Что качается темы дано статьи, т.е. алгоритмов поиска можно одназначно отметить одну вещь – CUDA эффективна в таких алгоритмах, где CPU вынужден повторять итерации, а CUDA может сделать то же самое в разных потоках. Другими словами, если при переносе алгоритма с CPU на GPU мы избавляемся от цикла и переносим его на потоки.
Так же нужно отметить общее высказывание, о том, что CUDA эффективна в таких задачах, где на входе имеем большой массив данных и на выходе так же имеем большой массив данных.

СПИСОК СОКРАЩЕНИЙ

CUDA - Compute Unified Device Architecture
GPU - Graphic Processor Unit
CPU - Central Processor Unit

СПИСОК ЛИТЕРАТУРЫ

1. Кнут Д.Э. Искусство программирования, том 3. Сортировка и поиск, 2-е изд. — М.: «Вильямс», 2007. — С. 824. — ISBN 0-201-89685-0
2. Parallel search on video cards, article, oracle server technologies – special projects

Автор: Бокатюк Олег

Административное: новые разделы форума

lexa — Sun, 18 Apr 2010 12:34:18 +0000

По многочисленным просьбам, заведены новые разделы Форума:

Старые топики будут частично перенесены в новые разделы в рабочем порядке (дело это ручное, а потому мучительное).

Кроме того, напоминаю о двух возможных удобствах, важных если вы размещаете исходный текст программы или его кусок:

Начало/конец фрагмента кода нужно помечать тегами <code>.... </code>: В этом случае парсер не съест знаки больше-меньше, разметит номера строк и читать ваш код будет куда приятнее.
Если хочется разместить большой фрагмент или целую программу и т.п.,: то при создании топика в форуме надлежит использовать закладку File attachments, тогда ваш архив будет прицеплен к топику. Это возможно только для стартового сообщения топика и только для зарегистрированных/авторизованных пользователей, но, увы, не для комментариев и не для анонимусов.

Кроме того, напоминаю, что я (как админ сайта) заинтересован в авторах. Если вы хотите писать не в форум, а "на первую страницу", пишите через contact form и вы будете повышены до автора.

Про CUDA на русском

lexa — Fri, 16 Apr 2010 08:35:13 +0000

Судя по регулярно задаваемым тут вопросам, эта тема многих интересует, поэтому анонсирую.

Довольно давно анонсированная книга по CUDA начала, наконец, продаваться:

дешевле всего в Озоне (294 рубля), но там дата поступления в доставку- 26 апреля. В books.ru - поступление ожидается 20-го, а Болеро просто пишет что отправят за день (но не проверял).

Из аннотации:

Данная книга посвящена программированию современных графических процессоров (GPU) на основе технологии CUDA от компании NVIDIA. В книге разбираются как сама технология CUDA, так и архитектура поддерживаемых GPU и вопросы оптимизации, включающие использование .PTX.
Рассматривается реализация целого класса алгоритмов и последовательностей на CUDA.
К книге прилагается CD, который содержит примеры решения на CUDA реальных задач с большим объемом вычислений из широкого класса областей, включая моделирование нейронных сетей, динамику движения элементарных частиц, геномные исследования и многое другое.

Не читал, поэтому пока про содержание не имею мнения.

P.S. Спасибо компании NVidia (московской) за мониторинг состояния издания книги.

OpenCL: ATI + NVidia, текущее состояние дел (windows)

lexa — Sat, 20 Feb 2010 13:56:42 +0000

По случаю выхода ATI Stream SDK 2.01, описанное в предыдущей статье несколько устарело. Теперь все иначе, по меньшей мере под Windows.

OpenCL.DLL устанавливается и драйверами NVidia (поверх имеющегося) и ATI SDK (только если в системе такого нет? если в системе есть более старый?) . При этом тот вариант, который ставят драйвера NVidia 196.xx - хороший, позволяющий работать и с NVidia и с ATI картами без каких-то избыточных телодвижений, а вот ATI-шный в двухкарточном варианте у меня вызывал непонятные проблемы.

Другими словами, надежнее после установки ATI SDK установить свежие драйвера NVidia еще раз.

Поддержка OpenCL ICD наконец то заработала, поэтому находятся следующие устройства:

Если к обоим картам подключены мониторы, а в пути (PATH) находится atiocl.dll (atiocl64.dll на 64-битной системе), то будут найдены оба GPU (+CPU-устройство, обеспечиваемое драйверами ATI).
Если от NVidia отключить монитор, то устройство пропадает из списка найденных (трюк с запуском PhysX на такой карте не работает, собственно такой настройки нет).
Если отключить монитор от ATI-карты, то Platform остается в списке устройств, однако судя по скорости расчета - он происходит на CPU.

Кроме того, интеграция с OpenGL полностью работает только если карта NVidia сделана primary display (скажем, в примере oclNbody отрисовка параметров работает во всех случаях, а отрисовка результатов расчетов - только на primary display).

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

Административное: новый поиск

lexa — Mon, 08 Feb 2010 19:21:58 +0000

Собирался года два, наконец собрался. Поиск на сайте заменен на наш («Ашманов и Партнеры»).

Теперь русский язык - склоняется (английский - тоже, но это и раньше было), если вы искали шейдер, то шейдеры тоже найдутся. Поддерживается язык запросов: кавычки, и-или-не (краткое описание есть тут).

Индексируются и тексты и комментарии.

Я ниасилил сделать модуль для Drupal, вместо этого асилил jQuery.

Должно работать везде (как jquery и положено), ну может кроме MSIE 5.x (и не жалко), при обнаружении каких-то проблем просьба писать в комментарии в этом сообщении.

P.S. Листалку делать тоже поленился. Выдаются первые 30 результатов, если их реально больше, то к 31-му результату никак не попасть, используйте "минус-слова" (чтобы поискать все тексты где есть OpenCL и нет NVidia нужно задать запрос OpenCL ~nvidia).

NVdia Nexus: работа из виртуальной машины

lexa — Thu, 04 Feb 2010 15:00:53 +0000

После успеха с запуском NVidia Nexus на двух машинах, я загорелся сделать это на одной.

Стандартные системные требования в этом случае простые: две G92/G200 видеокарты, к одной подключается монитор (мониторы, если их несколько), вторая настраивается как "ускоритель PhysX" (без подключенного монитора) и все должно работать.

У меня в рабочей машине - две видеокарты, но одна из них NVidia, а вторая - ATI и сделать так, как написано выше - невозможно, если отключить от NVidia-карты мониторы, то она пропадает из системы, запустить CUDA-программу даже локально (без Nexus) не выходит. Вынимать ATI (менять на NVidia) не хотелось, для третьей двухслотовой карты места в машине нет, покупать еще одну видеокарту хотелось еще меньше.

Первая попытка: локальный запуск (epic fail)

Если работать локально, то все работает.... до первой остановки на breakpoint внутри kernel. Даже если все активные окна расположены на мониторе, подключенном к ATI, в момент остановки на точке остановки все зависает.

Судя по особенностям процесса установки в локальном режиме ("рабочая" карта должна стать PhysX-ускорителем), в любых других режимах и должно виснуть.

Вторая попытка: виртуальная машина

В то же время, при удаленной отладке ATI-шный монитор остается живым и активным, а значит зависания - не обязательное свойство.

Завожу второй Windows в виртуальной машине (VMWare), ставлю туда Visual Studio, ставлю туда Host-часть Nexus, настраиваю удаленную работу......все работает, включая и пошаговую отладку. Естественно, окно виртуальной машины располагаю на ATI-шном мониторе.

Я не уверен, что все работает "100-% устойчиво", но пару десятков сессий пошаговой отладки удалось провести без малейших проблем.

Получается, что с виБрртуальной машиной (со второго компьютера) у меня в первый раз не завелось по причине ~~общей жизненной неустроенности~~ каких-то проблем с установкой Nexus Monitor, тогда же у меня вообще почти ничего не завелось, ни синхронизация, ни secure access.

Мораль

Требование двух G92/G200 видеокарт является необязательным, вполне можно обойтись одной G92/G200 (для запуска CUDA) и одной какой-нибудь.

Аппаратные подробности

Windows7/x64, Core i7, много дисков и памяти.
Primary Display (в терминах Windows) - подключен к ATI HD5870, драйвера ATI последние (10.1)
Второй монитор - подключен к GTX280, драйвера рекомендованные для Nexus (195.62)
На самой машине пускаем Nexus Monitor
В виртуальной машине (тоже Windows 7) пускаем Nexus Host, подключаемся к настоящей машине "по сети".
Окно виртуальной машины на primary display, поэтому неотрисовки и т.п. на втором мониторе (при висении в отладчике) - не волнуют.
Виртуальная машина - VMWare Workstation, 7.0.1, VMWare tools установлены (включая дисплейный адаптер)
Из возможно важных настроек (но менять не пробовал) - сеть в VMware bridged.

И еще

Отсутствие conditional breakpoints очень сильно уменьшает полезность. Если мне нужен не пошаговый режим, а именно остановка в случае "А меньше 5, Б больше 8", то придется такой if() рисовать, перекомпилировать и в нем и останавливаться. Прям девятнадцатый век...

NVidia Nexus: нулевые шаги

lexa — Mon, 01 Feb 2010 18:15:05 +0000

UPDATE: от большинства проблем удалось избавиться просто переустановкой NVidia Nexus на обеих машинах. Secure connection заработала, синхронизация - преимущественно заработала, хотя бывают необъяснимые сбои (и тогда надо удалить синхронизируемые файлы с той машины, на которую копируем т.е. ту, где Nexus Monitor).

Update2 И с виртуальной машиной тоже получилось, поэтому проще поменять полностью текст, чем вносить еще поправки:

Если не работает синхронизация или Secure connection (при стандартных настройках, когда вы ничего не меняли) - попробуйте переставить Nexus Monitor, мне помогло.
Иногда (закономерность понять не удалось) синхронизация файлов не может синхронизироваться. Тогда нужно вручную удалить синхронизируемые файлы с той машины, где запущен Nexus Monitor (на Win7/Vista это %user profile%/AppData/Roaming/Nvidia corporation/Nexus/Monitor/Mirror/)

Курс по CUDA для всех желающих (Москва)

lexa — Fri, 22 Jan 2010 14:28:23 +0000

В прошлом году я прощелкал, а в этом - нет, успеваю анонсировать.

С 24 февраля по 12 мая, еженедельно, по вторникам, на ВМиК МГУ (Москва, Воробьевы горы, м. Университет) будет читаться второй уже курс программирования NVidia CUDA для всех желающих.

Процитирую из анонса:

Московский Государственный Университет им. М.В. Ломоносова совместно с компанией NVIDIA приглашает заинтересованных студентов пройти специализированный курс "Архитектура и программирование массивно-параллельных вычислительных систем" на основе технологии CUDA. В рамках курса вы узнаете о современных многоядерных архитектурах, моделях программирования и основополагающих принципах, лежащих в основе построения эффективных параллельных алгоритмов. Вы также познакомитесь с реализациями типичных алгоритмов и задач, возникающих в цифровой обработке сигналов, математическом моделировании и гидродинамике. По окончании курса вы сможете применить свои знания на практике уже сегодня при решении вычислительноемких задач в ваших курсовых и дипломных работах. Приобретенные знания необходимы для всех, кто планирует связать свое будущее с высокими технологиями и высокопроизводительными вычислениями. Всем студентам, успешно завершившим курс "Архитектура и программирование массивно-параллельных вычислительных систем" будут выданы дипломы.

Ссылки:

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

Еще об NVidia и ATI в одной машине

lexa — Thu, 31 Dec 2009 14:01:23 +0000

Как ни печально, но поведение драйверов NVidia ничем не лучше драйверов ATI: чтобы CUDA/OpenCL работали, к какой-то видеокарте NVidia должен быть подключен какой-то из активных дисплеев.

Таким образом, идея потестировать NVidia Nexus на одной машине с двумя видеокартами, одна ATI и одна NVidia, полностью провалилась: если к Geforce не подключать монитор, то с точки зрения CUDA/NEXUS этой карты не существует. Вот если бы у меня было две карты NVidia, разговор был бы другой. Еще одна карта NVidia у меня есть, но слоты кончились....

Конечно, у меня есть еще немножко компьютеров и Nexus я найду как потестировать (или вытащу ATI, вставлю G80), но обидно, да. Ограничение чисто софтверное и явно на пустом месте (ATI это тоже касается, та же фигня).

AMD/ATI Stream SDK 2.0

lexa — Wed, 23 Dec 2009 19:03:17 +0000

Вышел ATI Stream SDK 2.0 (релиз).

Лично меня интересует только OpenCL, поэтому и смотрел я исключительно с этих позиций:

Примеры, ранее требовавшие primary display на карте ATI (NBody) - теперь стали работать нормально, главное OpenCL.dll подсунуть от ATI.
Появилась интеграция OpenCL-OpenGL, но бинарной совместимости с таковой интеграцией у NVidia нет, поэтому примеры из NVIdia GPU Computing SDK без перекомпиляции на ATI не работают (а с перекомпиляцией - не пробовал). Мне не очень нужно, но из общих соображений кажется, что починят.
Есть поддержка OpenCL ICD (Installable client drivers). В теории это должно дать возможность выбирать аппаратуру из имеющейся на ходу. По факту - этого не происходит, во всяком случае oclDeviceQuery/CLInfo показывают только ту карту, OpenCL.dll от которой попался первым в пути.
В этом топике форума AMD пишут, что возможно OpenCL.dll разных версий (размер - точно разный), таким образом есть надежда, что проблема как-то будет починена.
Появился профайлер, хотя с первого взгляда не очень понятный.

Больше ничего толком за полдня не посмотрел, но общее впечатление такое, что с beta4 стало несколько лучше.

ATI HD4xxx и OpenCL

lexa — Sun, 20 Dec 2009 16:26:48 +0000

Читал форумы AMD, много думал.

Вычитал грустное про HD4xxx и OpenCL:

The HD4XXX series was not designed for OpenCL whereas the HD5XXX series was. There will be performance improvements on this series because of improvements in the HD5XXX series, so it will get better, but it is not our focus.

For example, if you are using local memory, they are all currently emulated in global memory. So it is possible you are going out to main memory twice as often as you do on NVidia. This can cause a fairly large performance hit if the application is memory bound. On the HD5XXX series, local memory is mapped to hardware local and thus is many times faster than the HD4XXX series.

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

Другими словами, если вас интересует результат, то под 4xxx надо программировать сильно иначе, например локальный storage делать на регистрах (если их хватит).

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

OpenCL, NVidia, ATI и все все все....

lexa — Sun, 20 Dec 2009 16:11:26 +0000

Вопрос о совместимости OpenCL между железом разных производителей - является необычайно волнительным. Понятно, что задумывалось все совместимым, но насколько это получилось и сколько граблей разложено на пути к этой совместимости - отдельный интересный вопрос, который перешел в практическую плоскость сразу после выпуска ATI Stream SDK 2.0 beta 4, где появилась поддержка OpenCL на графическом адаптере. До того поддержка OpenCL на PC была только у NVidia и вопрос был сугубо теоретическим.

Как только мне в руки попал Radeon HD5870, так сразу захотелось тему исследовать... и в первом приближении я ее изучил.

Первое приближение: жизнь есть

Более того, какая-то совместимость достигается и без перекомпиляции приложений. Если быть точным то:

Часть примеров из NVidia SDK работает на оборудовании ATI, часть - работает частично (т.к. нет нужных свойств, например интеграции OpenCL/OpenGL), часть - не работает. Все проверялось без перекомпиляции/перелинковки приложений, на готовых бинарных файлов.
Аналогично, часть примеров ATI SDK работает без проблем, часть - ругается, дескать не хватает чго-то.

Но в первом приближении - жить можно. Да, надо иметь зоопарк железа для тестирования т.к. архитектура везде разная и оптимальные kernels, соответственно, будут тоже разными.

DLL Hell

DLL Hell имеется в полный рост, причем есть все предпосылки к тому, что дальше будет еще хуже:

OpenCL.dll для карты NVidia поставляется вместе с драйвером и впоследствии оказывается в \Windows\SysWOW64 (у меня 64-битная система, на 32-битной наверное окажется в Windows\System32)
OpenCL.dll для ATI - является частью SDK и ставится в \Program Files\ATI Stream

Соответственно, в зависимости от $PATH приложение работает с одной, либо другой видеокартой.

Как только ATI включит эту DLL в поставку драйверов - сразу начнется драка с NVidia в полный рост, я предсказываю это.

Странности с ATI

Часть примеров ATI SDK работает нормально только если монитор, подключенный к Radeon является primary display. Код этих примеров я внимательно не рассматривал, на невнимательный взгляд такого быть не должно бы. Нужно разбираться, возможно это происки бета-версии.

Производительность

Оценка производительности - задача сложная, особенно с учетом разницы в архитектурах. Поэтому приведу лишь одну цифру, относиться к которой нужно крайне осторожно. Один шаг примера Nbody (из ATI SDK) с числом частиц 9984 исполняется, согласно его внутренней измерялке:

280 миллисекунд на CPU (i7, 3.6Ghz)
28 миллисекунд на ATI HD5870
12 миллисекунд на NVidia GTX280

При увеличении размеров задачи относительная разница между GPU уменьшается (66/45 ms для 19968 частиц), возможно дело вообще в стадии инициализации или еще в чем-то подобном. В то же время, я ожидал, что HD5870 будет заметно быстрее GTX280, возможно архитектура NVidia больше подходит именно к задаче многих тел.