CUDA: аспекты производительности при решении типичных задач

    Перед тем как начать переносить реализацию вычислительного алгоритма на видеокарту стоит задуматься — получим ли мы желаемый прирост производительности или только потеряем время. И несмотря на обещания производителей о сотнях GFLOPS, у современного поколения карт есть свои проблемы, о которых лучше знать заранее. Я не буду глубоко уходить в теорию и рассмотрю несколько существенных практических моментов и сформулирую некоторые полезные выводы.

    Будем считать, что вы примерно разобрались, как работает CUDA и уже скачали стабильную версию CUDA Toolkit.

    Я буду мучить теперь уже middle-end видеокарту GTX460 на Core Duo E8400.

    Вызов функции

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

    __global__ void stubCUDA(unsigned short* output)
    {
      // the most valid function: yep, does nothing.
    }

    Напомню, что спецификатор __global__ позволяет выполнить функцию на GPU, вызвав ее с CPU:
    cudaThreadSynchronize();
    stubCUDA<<<GRID, THREADS>>>(0);
    cudaThreadSynchronize();

    Все вызовы функций по умолчанию асинхронны, поэтому вызовы cudaThreadSynchronize() необходимы для ожидания завершения вызванной функции.

    Попробуем прогнать такой блок в цикле: получаем порядка 15000 вызовов в секунду для GRID=160, THREADS=96.

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

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

    Проверим. Без синхронизации удалось запустить функцию 73100 раз в секунду. Результат, надо заметить, нисколько не впечатляющий.

    И последний тест, запустим функцию с GRID=THREADS=1, казалось бы, это должно устранить накладные расходы на создание кучи потоков внутри карточки. Но это не так, получаем те же 73000-73500 вызовов в секунду.

    Итак, мораль:
    • Абсолютно бессмысленно запускать на карточке те задачи, которые и на CPU считаются за миллисекунды.
    • Синхронизация потоков после вызова уменьшает производительность совсем незначительно на средних задачах.
    • Количество потоков и размер сетки не влияет на итоговое количество вызовов в секунду (разумеется это не так для «полезных» функций, которые делают что-то).

    Доступ к памяти извне

    Для того чтобы считать что-то полезное нам потребуются входные и выходные данные. Для этого надо понимать насколько быстро идет передача данных из/в видеокарту. Воспользуемся следующей функцией:
    cudaMemcpy(data_cuda, image, data_cuda_size, cudaMemcpyHostToDevice);

    Да, CUDA предлагает нам и средства асинхронной передачи данных, но их производительность, забегая вперед, не отличается от синхронной функции.

    Копируем большие блоки: как и в сторону cudaMemcpyHostToDevice, так и cudaMemcpyDeviceToHost получаем производительность порядка 2 Гбайт/c на больших блоках (более 100 мегабайт). В целом это очень даже неплохо.

    Значительно хуже обстоят дела с совсем небольшими структурами. Передавая по 4 байта мы получаем не более 22000 вызовов в секунду, т.е. 88 кбайт/c.

    Мораль:
    • Желательно группировать данные в большие блоки и передавать их одним вызовом функции cudaMemcpy.

    Доступ к памяти изнутри

    После того как мы передали данные на карточку, можно начинать с ними работать. Хочется оценить примерную скорость доступа к видеопамяти. Для этого напишем следующую функцию:
    __global__ void accessTestCUDA(unsigned short* output, unsigned short* data, int blockcount, int blocksize)
    {
      // just for test of max access speed: does nothing useful
      unsigned short temp;
      for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
      {
        int vectorBase = i * blocksize;
        int vectorEnd = vectorBase + blocksize;

        for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
        {
          temp = data[j];      
        }
      }
      output[0] = temp;
    }


    Здесь уже используются параметры GRID и THREADS, пока не буду объяснять зачем, но поверьте — все как следует. Придирчивые скажут, что результат пишется неправильно из-за отсутствия синхронизации, но нам-то он и не нужен.

    Итак, получаем порядка 42 Гбайт/c для произвольного чтения. Вот это совсем неплохо.

    Теперь модифицируем функцию, чтобы она копировала входные данные на выход. Бессмысленно, но позволяет оценить скорость записи в видеопамять (поскольку изменение совсем несложное, я не буду дублировать код).

    Получаем порядка 30 Гбайт/с на ввод-вывод. Тоже неплохо.

    Следует сделать поправку на то, что фактически мы использовали последовательный (с некоторыми отступлениями) доступ к памяти. Для произвольного цифры могут ухудшится до двух раз — но ведь и это не проблема?

    Мораль:
    • За счет очень высокой скорости доступа к памяти на картах эффективно реализовывать алгоритмы, интенсивно ее использующие.

    Арифметические операции

    Совсем простые примеры опустим и сделаем что-то полезное. А именно — нормализацию изображения (pixel[t] := (pixel[t]-sub)*factor). Собственно код:
    __global__ void normalizeCUDA(unsigned short* data, int blockcount, int blocksize, float sub, float factor)
    {
      for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
      {
        int vectorBase = i * blocksize;
        int vectorEnd = vectorBase + blocksize;

        for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
        {
          register float d = (float)data[j];
          d = (d - sub) * factor;
          data[j] = (unsigned short)d;
        }
      }
    }


    Здесь используется аж три казалось бы затратных вычислительных процедуры: приведение к вещественным числам, ADDMUL и приведение к целым. На форумах пугают, что приведение целые-вещественные работает из рук вон плохо. Может быть это было верно для старых поколений карточек, но сейчас это не так.

    Итоговая скорость обработки: 26 Гбайт/c. Три операции ухудшили производительность относительно прямого ввода-вывода всего на 13%.

    Если внимательно посмотреть код, то нормализует он не совсем верно. Перед записью в целые числа, вещественное необходимо округлить, например функцией round(). Но не делайте так, и постарайтесь ее никогда не использовать!

    round(d): 20 Гбайт/c, еще минус 23%.
    (unsigned short)(d + 0.5): 26 Гбайт/с, собственно время в пределах погрешности измерений даже не поменялось.

    Мораль:
    • Арифметические операции работают и правда быстро!
    • Для простейших алгоритмов обработки изображений можно рассчитывать на скорость в 10-20 Гбайт/c.
    • Лучше избегать использования функции round().

    Логические операции

    Попробуем оценить скорость работы логических операций, а заодно сделаем еще одно доброе дело: найдем минимум и максимум значений в массиве. Данный этап обычно предваряет нормализацию (и именно для этого и писался), но у нас все будет наоборот — т.к. он сложнее. Вот рабочий код:
    __global__ void getMinMaxCUDA(unsigned short* output, unsigned short* data, int blockcount, int blocksize)
    {
      __shared__ unsigned short sMins[MAX_THREADS];
      __shared__ unsigned short sMaxs[MAX_THREADS];
      
      sMins[threadIdx.x] = data[0];
      sMaxs[threadIdx.x] = data[0];
      
      for (int i = blockIdx.x; i < blockcount; i += gridDim.x)
      {
        int vectorBase = i * blocksize;
        int vectorEnd = vectorBase + blocksize;

        for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
        {
          register unsigned short d = data[j];
          if (d < sMins[threadIdx.x])
            sMins[threadIdx.x] = d;
          if (d > sMaxs[threadIdx.x])
            sMaxs[threadIdx.x] = d;    
        }
      }

      __syncthreads();

      if (threadIdx.x == 0)
      {
        register unsigned short min = sMins[0];
        for (int j = 1; j < blockDim.x; j++)
          if (sMins[j] < min)
            min = sMins[j];
        if (min < output[0])
          output[0] = min;
      }
      
      if (threadIdx.x == 1)
      {
        register unsigned short max = sMaxs[0];
        for (int j = 1; j < blockDim.x; j++)
          if (sMaxs[j] > max)
            max = sMaxs[j];
        if (max > output[1])
          output[1] = max;
      }

      __syncthreads();
    }


    Здесь уже не обойтись без синхронизации потоков и shared memory.

    Итоговая скорость: 29 Гбайт/c, даже быстрее нормализации.

    Почему я объединил код минимума и максимума — обычно нужны оба, а вызовы по отдельности теряют время (см. первый абзац).

    В общем, киньте камнем в того, кто сказал что на видеокартах плохо с условным операциями: искусственно удалось замедлить этот фрагмент практически в 2 раза, но для этого потребовалось увеличить глубину условий аж до 4! if () if () if () if () else if ()…

    Мораль:
    • На современных карточках в общем-то не так и плохо с логикой, но следует избегать большой глубины вложенных условий.

    Сложные структуры данных

    Руководствуясь идеей о том что алгоритмы и структуры данных сильно связаны (хотя бы вспомнить Н. Вирта), следует проверить как же обстоят дела с некоторыми сложными структурами данных.

    Вот тут возникает проблема, при передаче данных в функции мы можем использовать только два вида объектов — константные интегральные типы (чиселки) и ссылки на блоки видеопамяти.

    Идея строить например деревья, основанные на ссылках накрывается сразу же:
    • мы не можем выделять память из функции, работающей на карточке;
    • любое выделение и копирование небольшого объема данных работает очень медленно (см. раздел 2).

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

    Расплата за подобные ухищрения — необходимость применения двойной индексации:
        for (int j = vectorBase + threadIdx.x; j < vectorEnd; j += blockDim.x)
        {
          temp = data[index[j]+i];      
        }

    Данный фрагмент работает со скоростью от 10 до 30 Гбайт/c в зависимости от наполнения и размеров индекса и данных. Использование памяти можно пытаться соптимизировать но даже в лучшем случае мы теряем 25% скорости доступа. Тройные индексы ведут себя еще хуже, теряя 40%-60% производительности.

    Сегодня мы многое поняли

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

    Если топик вам понравится, то расскажу про то как обсчитывать на видеокарте несколько полезных объектов: Distance Map, морфологию изображений и поисковые индексы и покажу несколько интересных структур данных, которые работают достаточно быстро и не создают лишних проблем с синхронизацией.
    Share post

    Comments 67

      +1
      Ну что тут сказать… В избранное!
      +1
      Если топик вам понравится, то расскажу про то как…
      Топик нам понравился. Хочется действительно больше реальных применений, которые можно использовать в своих проектах. Пока что только кодирование аудио/видео припоминается.
        0
        Обработка изображений, фурье-анализ (хотя это туда же), решение систем уравнений, моделирование методом конечных элементов.
          0
          До учёных эта штука пока не очень докатилась, как докатится — начнём считать всевозможные коллапсы и прочее.
            0
            Я видел применения CUDA в научных работах с разных факультетов нашего университета. На факультете, где учусь я, основная проблема — отсутствие программистов. Я даже знаю лаборатории, где есть компьютеры с мощными карточками, но нету людей, которые с ними работают.

            В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
              0
              > Я видел применения CUDA в научных работах с разных факультетов нашего университета.
              Список публикаций, пожалуйста.

              Даже больше скажу: scholar по запросу CUDA выдаёт только статьи с таким учёным, по запросу же OpenCL выдаёт статью-спецификацию OpenCL на которую уже есть 21 ссылка (за ~1 год)

              > В Дубне знаю людей, которые уже давно с помощью CUDA циклотроны считают.
              И что-же они считают и как? Какие именно «циклотроны»?
                0
                Отвечу сам себе: FAIR использует CUDA: mpd.jinr.ru/trac/browser/trunk/cuda

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

                  По Дубне можно найти тут информацию groups.google.com/group/cudacsmsusu/?hl=ru. Я с телефона, мне трудно найти ссылку сразу до нужного файла. Там где-то в прикрепленных файлах должна быть презентация.
                    0
                    Я нашёл уже, спасибо. Будем посмотреть.
                    0
                    И ещё раз — нет, не только там:
                    www.iki.rssi.ru/seminar/20100225/perepelkin.pdf

                    Модель там, конечно, упрощённая. Но за наводку спасибо.
                    0
                    Хотя вру — по CUDA есть две статьи, суммарное число ссылок порядка 500.
                0
                  0
                  А как насчет рендеринга в 3DS Max? Можно ли CUDA приспособить под это и какой будет прок?
                  0
                  Спасибо. Сам полгода назад бодался с технологией на реальном коде и писал по итогам статью habrahabr.ru/blogs/hi/116579/
                  Были большие проблемы с ветвлениями, возможно последняя CUDA оптимизирована лучше. Впрочем, карта была послабее — gts450
                  А не мерили сколько насколько полученная производительность ниже теоретической?
                  И что профайлер на if ы показывает?
                    +1
                    Профайлер все меряет на уровне функций. Если вы про число ветвлений, то их будет 0, поскольку столь маленькие if'ы считаются с помощью предикатов.
                      0
                      Ой, а можно насчет предикатов подробнее?
                        +1
                        Есть несколько (по-моему 4) однобитных предикатных регистров. Соответственно команды сравнения могут записывать в них результат, а затем для любой другой команды можно указать, чтобы она исполнялась только если предикатный регистр установлен в false или true. Так что в любом случае исполняются обе ветки в if'e, но не создается нагрузки на шедулер (которому нужно в начале разбить варп, а потом как-то соединить).
                        0
                        + некоторые if'ы сводятся к арифметическим выражениям. нет уверенности, что это именно так, но попытки их как-то оптимизировать вручную дали тот же результат что и исходная версия.
                          +1
                          Минимум и максимум реализованы аппаратно. ?: скорее всего тоже.
                      0
                      Спасибо, полезная статья.
                        +1
                        И еще, насчет условий. Сколько kernelов запускалось параллельно?
                        Насколько я помню, любой if рубит производительность на 2 * время выполняемого под ним кода (сначала выполнятся одновременно kernelы на TRUE, потом на FALSE
                        Соответственно, ifы меньше влияют если
                        1) Под ними простая арифметика. Попробуйте поставить под условие что-то посерьезней одного действия из быстрой памяти
                        2) GPU недогружен, т.е. количество одновременно выполняемых kernelов меньше количества исполнимых блоков на карте.
                          0
                          пробовал разные конфигурации, 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.
                            +1
                            Однако, скорость в ваших примерах достигается в основном не за счет параллельности выполнения, а за счет быстрой памяти…
                            Иначе бы от количества одновременно использующихся вычислительных блоков зависела бы почти линейно.

                              0
                              Верно. Скажем так, за счет быстрой памяти можно эффективно обрабатывать очень большие объемы данных, при этом стараясь сделать так, чтобы обсчет не сильно испортил это время.

                              Про именно вычислительные задачи напишу потом, однако там не все так радужно.
                                0
                                Классическая хорошо ускоряемая на граф. ускорителе задача — молекулярная динамика.
                                Когда модель загружается в память ускорителя и там «трясется» сутками, довольно сложной математикой.
                                Я в основном с такими задачами встречался, и ваш пример очень интересен т.к. он в противоположен.
                                  0
                                  интересно, а какое приблизительно получилось получить ускорение относительно CPU?

                                  в задачах, которые мы решаем в основном очень большие объемы данных фигурируют, и единственная цель, написать их обработку, так чтобы время передачи на карточку было соизмеримо со временем работы алгоритма, так что конечно все эти 10-40 Гбайт/с больше, чем нужно. хотя многие алгоритмы работают в несколько проходов — там это важно.
                                    +1
                                    Ну, до самостоятельного написания молекулярной динамики я не дорос ещё. Но NVIDIA хвастается серьезным ускорением www.nvidia.com/object/tesla_bio_workbench.html
                                    Через месяц будет Tesla, проверим на практике.

                                    А в моей задаче (сложный ассоциативный поиск по большой базе данных) удалось сравниться с i5 на OpenMP. Но это, конечно, не тот результат, которого хотелось.
                                    Там еще, кроме ветвлений, проблема конкурентного доступа к памяти на чтение.
                                      0
                                      о, спасибо. как ни странно не сталкивался с готовой библиотекой (хотя, правда, она нам не очень поможет, но интересно).

                                      несколько не понятно — а в чем проблема, если производится только чтение… или имелась в виду проблема с чтением/записью?
                                        0
                                        Там блокировки по чтению одной ячейки несколькими потоками. Увы. В Fermi эта проблема менее остро стоит, но стоит. А в константную память не получается засунуть.
                                          0
                                          не зная конкретной задачи могу только порассуждать о правильном использовании адресации, чтобы подзадачи пересекались по-минимуму.

                                          есть такая интересная особенность, что потоки решающие одинаковую по размеру подзадачу обычно завершаются одновременно, этот факт можно косвенно использовать (например один идет с начала блока — второй с конца).
                                            0
                                            и кстати, в процессе решения одной задачи мы придумали забавный метод — наплевать на синхронизацию вовсе, а после вычисления свалидировать значения и где нужно подправить — конечно, специфика это.
                                0
                                И все-таки, если тест под рукой — попробуйте засунуть под пару ifов что-то поболее, например цикл.
                                Очень интересны результаты.
                              0
                              >Количество потоков и размер сетки не влияет на итоговое количество вызовов в секунду
                              Ой-ой-ой.
                              Советую попробовать запустить пустой кернел в конфигурации GRID=1024, THREADS=512 и посчитать количество запусков в секунду. (будет сильно меньше)
                              Если надо, могу объяснить почему.
                                0
                                запустил (1024,512); те же 71000 получается. зачем мне врать? )
                                  0
                                  да, после того как написал коммент, понял что ступил. извиняюсь.
                                0
                                еще дополнение вместо передачи blockcount и blocksize в код кернела, можно использовать build-in переменные dim3 blockDim и dim3 gridDim
                                  0
                                  в общем-то они и передаются (и используются, как же без них), но они не обязаны быть кратными размерности решаемой задачи, поэтому и передаю дополнительные параметры, которые можно как-то покрутить. в общем да, ничего серьезного, производительность отличается на единицы процентов.
                                  0
                                  когда уже будут процессоры синхронизировать с видеокартами, т.е. потоки сами бы распределялись в зависимости от задачи/сложности )))
                                    0
                                    в общем-то спецификация __global__ позволяет выполнять одни и те же функции на CPU и GPU. но смысла в этом не очень много, мощностями CPU можно просто пренебречь.

                                    а вот переписать код, который работает в рамках классической модели на видеокарты — задача, которая автоматически (пока) не разрешима — другая модель, все-таки.
                                      0
                                      Читал как-то в журнале Суперкомпьютеры, что есть проект поставить кластер под OpenCL.
                                      Опять же, Intel уже выпустил альфу OpenCL для своих процессоров.
                                      Было бы очень интересно.
                                      А мощностями современных CPU пренебрегать не стоит, начиная с Nehalem они начали сокращать отставание с видеокартами.
                                        0
                                        интересна примерная производительность i5/i7 в GFLOPS.

                                        прогресс по обоим направлениям идет, однако пока видеокарты устойчиво выигрывают в отношении производительность/стоимость и производительность/энергопотребление. хотя, конечно, это может зависеть от задачи, в частности задачи, которые вообще не параллелятся, а одно ядро CPU все-таки значительно мощнее ядра GPU.
                                          0
                                          Производительность CPU «отстаёт» от GPU. У CPU сейчас Тера-Флопы а у CPU «всего» Гига-Флопы. 5 Тера-Флопс (AMD 6990) vs 128 Гига-Флопс (Core i7). Проблем у GPU только две: PCI-Express и алгоритм. Не каждый алгоритм можно посадить на графическую карту. И одним из главных «бутылочных горлышек» это PCI-Express на графику и обратно в процессор. OpenCL хорош в том, что может использовать оба два типа устройств и хорош в том, что можно использовать «совместную» память как для CPU и GPU, но кернелы должны бы быть по идее разные для каждого типа.
                                        0
                                        Только не __global__. __global__ означает функцию, которая выполняется на GPU, а вызывается с CPU.
                                        Если нужны обе версии функции, то нужно писать сразу 2 спецификатора: __device__ __host__. Компилятор сгенерирует две бинарные версии функции: одна будет вызываться и выполняться на CPU (этим правда займётся gcc/icc/cl), а другая делать тоже самое на GPU.
                                          0
                                          да, конечно! и кто меня попутал.
                                      +2
                                      Хотел бы немного уточнить то, как можно по-настоящему замедлить CUDA-программу.
                                      То, что ты называешь «случайным» чтением, таким не является. Как раз-таки наоборот, в твоём примере самый быстрый вариант доступа – coalesced (когда все нити одного варпа обращаются к одному блоку памяти). Он в 8-16 раз быстрее по-настоящему случайного доступа. Вообще доступ к памяти (как на чтение, так и на запись) – самая медленная операция на GPU (200 и больше тактов). Другое дело, что если блоков много, то пока одни ждут – другие считают.
                                      Нужно замедлить арифметические операции? Есть же остаток от деления! Вот где смерть производительности.
                                      Логические операции? Сделай так, чтобы все нити варпа бежали по разным веткам условия ( if… else if… else if ...). Вот тогда просмотрим.
                                      Справедливости ради хочется отметить, что CUDA CC 2.х устройства могут выделять память из пула с помощью malloc'ов…
                                      Спасибо за статью! Почитать про структуры данных было бы интересно.
                                        0
                                        Причём, я так понимаю, для чтения поток переводят на другой процессор.

                                        А всевозможные if/else/for сейчас исполняются не очень оптимально, по крайней мере в ATI.
                                          0
                                          Чтение происходит на другом процессоре, если оно идёт из текстур. Там хитрая схема преобразования адресов и данных. Если же программа обратилась к обычному адресу в RAM, то варп засыпает, происходит группировка обращений по соседним адресам и когда результат пришёл — скедулер снова переключается на этот варп.

                                          Не знаю как в AMD, но в CUDA с ветвлениями всё двояко. Если правильно распределить данные по потокам, то ветвления не навредят. Циклы тоже работают неплохо, хотя если их разворачивать, то это даёт небольшой прирост производительности.
                                            0
                                            Это вы по coalesced говорите, я правильно понимаю? Но ведь не все обращения такие?
                                              0
                                              Это общая схема. Если доступ к данным невозможно сгруппировать, то обращения полуварпа сериализуются и все эти потоки засыпают до тех пор, пока не приедут все результаты. То есть, если все 8 потоков полуварпа обратились к адресам, лежащим в 8 разных блоках памяти (1 блок – это выровненные по адресам 32/64/128 байтные сегменты в RAM), то делается 8 последовательных запросов к памяти.
                                        0
                                        Всё это хорошо, но где реальные программы использующие CUDA? Ну ладно, есть Премьер который использует ее. Всякие мутные шареваре-поделия типа МедиаЭспрессо и vReveal… А где массовые кодировщики видео/аудио/фото? Ниша практически пустая. Когда прикрутят к Megui, например?
                                          0
                                          CUDA – молодая технология. Далеко не все крупные компании готовы поставить на неё. Мало ли как завтра nVidia изменит API? Или вообще похоронит GPGPU на своих девайсах… Привязка к одному производителю видеокарт тоже играет роль.
                                          Но некоторые таки внедряют. Matlab, Photoshop да и многие другие пакеты, где SIMD архитектура походит, используют видеокарты (хороший пример – перебор паролей).
                                          Существуют реализации кодеков с использованием CUDA, но новые разработки с трудом опережают заоптимизированные насквозь реализации на CPU (с использованием SSE-инструкций и алгоритмических ухищрений).
                                          Реальные примеры использования CUDA можно увидеть в научной среде. Я вот некоторые куски своих прог так и ускоряю.
                                          Много примеров использования GPGPU можно увидеть в CUDA Zone
                                            0
                                            Корпорации да, но опенсорс мог бы и пошевелиться)
                                              +1
                                              В open source всё сложно. Есть открытые реализации некоторых интересных вещей. Например HOOMD для молекулярной динамики. Нужны хорошо параллелящиеся задачи, опытные алгоритмисты и достаточно много времени, чтобы разобраться в хитросплетениях оптимизаций для CUDA.
                                            0
                                            Есть кроссплатформенный opencl (ATI, Cell, NVidia), поэтому к CUDA никто цеплятся не хочет.
                                            0
                                            В университете изучаем возможности применения CUDA для решения плохо обусловленных СЛАУ. Встала проблема: точность.
                                            Двойная по стандарту IEEE754 поддерживается только с Compute Capability 1.3, и это пол-беды. Сам IEEE754 точность результата не описывает, а только точность операндов. Таким образом, берём плохо обусловленную СЛАУ, решаем в double extended числах, и получаем всего 17 знаков точности вместо «теоретических» 30+.
                                            Все матпакеты валятся на наших тестах, и реализации BLAS на CUDA — не исключение. С производительностью — да, хорошо, а вот с точностью — хуже, чем на CPU.
                                            Планируем писать библиотеку, строго обеспечивающую заданную точность элементарных операций. Пока в зачаточном состоянии, но, полагаю, можно будет «всего лишь» реализовать вычисления с произвольной точностью на CUDA.
                                              –1
                                              Попробуйте, для начала, переписать это на opencl.
                                                +1
                                                Чем это будет лучше? Для nVidia, OpenCL – обёртка над CUDA. Причём IMHO более слабая в плане API, работающая медленнее и при этом менее удобная.
                                                  0
                                                  Таки плюс у неё в портабельности.
                                                  Работает медленней она процентов на 5-10 (читал где-то).

                                                  Стоит поменять карточку, посмотреть что будет на других картах.
                                                    +1
                                                    Выигрыш может составлять от 10 до 60 %… Как повезёт.
                                                    Писать на CUDA удобнее и быстрее – она более развесистая, ближе к плюсам (есть даже шаблоны и какой-никакой STL). OpenCL – гораздо более низкоуровневая (спасибо, AMD) и урезанная технология. Для меня скорость разработки и расчётов важнее «портабельности». Проще в требованиях к железу указать какая карточка нужна. Да, я не разрабатываю клиентского ПО и не занимаюсь GRID-вычислениями, так что могу себе позволить такую блажь)
                                                    Со временем, разрыв между технологиями уменьшится или вообще исчезнет. Для OpenCL появится и нормальный API и продвинутые библиотеки и хорошая литература. Но я как-то дожидаться этого не хочу. Лучше я уже сейчас буду пользоваться преимуществами нового железа.
                                                      0
                                                      OpenCL — разработка Apple Inc. и Kronos Group. Или я чего-то не знаю?

                                                      А чем вы занимаетесь, если не секрет?
                                                        0
                                                        В 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 года, преподаю даже — рассказывааю студентам про С и Куду.
                                                          0
                                                          Мне почему-то кажется, что самый жёсткий вариант это не AMD даже, а pthreads/win threads.

                                                          Да, OpenCL v API 1.1.
                                                            0
                                                            Ну значит за последние полтора года ничего сильно не поменялось).
                                                            Дело в том, что CPU-потоки не накладывают практически никаких ограничений на распараллеливаемые куски кода. А вот GPU жёстко закрепляют список допустимых типов аргументов функций, делает всякие деления памяти на shared/global и прочие интересные вещи типа ограничений на размерности блоков/гридов. nVidia тщательно старается это всё скрыть и сгладить (иногда удачно, иногда не очень). Сейчас даже ООП можно применять на CUDA (если это кому-то надо)), серьёзно проработаны самые тормозные моменты, чтобы типичные для CPU операции не тормозили так сильно на GPU.
                                                            Кроме того CPU-потоки это полновесные процессы, которые делают сложные операции и практически не влияют друг на друга (ибо MIMD). Так как GPU – это SIMD, то нужно запускать потоки даже не по количеству процессоров, а с запасом в пару порядков и близко бегущие потоки будут сильно взаимодействовать. Поэтому я вообще слабо понимаю мотивацию для затраты кучи сил и времени на разработку такой технологии, с помощью которой можно работать на CPU также неудобно как на GPU). Лучше бы допилили GPGPU технологии.
                                                0
                                                К сожалению, в графических задачах точность не особо важна, поэтому с ней на GPU пока туго.
                                                Улучшить это дело – задача интересная и почётная. Только думаю придётся очень серьёзно покопаться в PTX, чтобы писать действительно быстрый код с большим количеством побитовых операций.
                                                0
                                                Код reduction (min/max) похож на код с сайта AMD, переписанный для CUDA ;-)
                                                  0
                                                  Книжка сильно в тему:
                                                  www.twirpx.com/file/413014/

                                                  Only users with full accounts can post comments. Log in, please.