Как реализовать поэлементное перемножение массивов средствами OpenMP, CUDA и GLSL?

Question

[Robotex]

Вот уже несколько дней бьюсь над решением этой проблемы, но мои решения получаются медленнее чем последовательное однопоточное решение:

for(int i=0;i<LAYERS;i++)
  {
    for(int j=0;j<INPUTS*NEURONS;j++)
    {
      temp[j] = inputs[j%INPUTS] * weights[i][j];
    }
    
    for(int j=0;j<NEURONS;j++)
    {
      inputs[j] = 0;
      
      for(int l=0;l<INPUTS;l++)
      {
        inputs[j] += temp[j*INPUTS + l];
      }
      
      inputs[j] = sigmoid(inputs[j]);
    }
  }

Т.е. если вкратце, то имеем массив в N элементов, массив в N*M, и массив (или матрица, здесь weight) в N*M*K элементов.

Поэлементно перемножаем первый массив на группу из K элементов третьего (когда в первом кончаются элементы, переходим на первый и вновь идем по его элементам) и результат во второй. Далее суммируем во втором массиве группы по M элементов и пишем суммы в первый массив (предварительно применив некоторую функцию, но не суть важно). Далее повторяем все с новой группой элементов третьего массива.


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

Answer 1

[svetlov]

Кажется, во втором варианте второй цикл не параллелится. Замените input[j] на локальную переменную, чтобы она в регистр упала. Еще может потребоваться промежуточный массив для складывания этих input[j] — чтобы потом одним махом переключить на него после выхода за цикл.

Answer 2

[Robotex]

Вот для CUDA пробую использовать cublas, перемножение матриц нашел, а вот как перемножить два массива не пойму.

Answer 3

[Robotex]

В чем принципиальная разница между этим: pastebin.com/8dpNKBwt (работает быстрее однопотокового аналога) и этим pastebin.com/fbe4gZSn (работает медленнее)?

Answer 4

[Robotex]

Так, с OpenMP решил вот так: pastebin.com/sJ4fXiAb Теперь вроде заметно ускорение.
Сейчас думаю над CUDA. пока тоже все получается вроде.

А вот с GLSL беда. Я вообще не знаком с этой технологией, может кто помочь?

Comments

[svetlov]

GLSL, CUDA или OpenCL? GPU одно, а технологии — разные. Хоть и похожие.

[Robotex]

OpenCL не нужно. У меня цель опробовать три технологии и выбрать ту, что будем использовать. Список тож не я составлял.

[svetlov]

Тогда под GLSL что имеется в виду? Это же всего лишь шейдерный язык для OpenGL.

[Robotex]

Да, шейдерный язык, но под GPGPU используется хорошо (и работает там, где недоступны CUDA и OpenMP).

И мне вот на этом шейдерном языке нужно максимально быстро и эффективно реализовать поэлементное умножение массивов.

[Robotex]

Например, видеокарты Intel не поддерживают ни CUDA, ни openCL, а если проц одноядерный, то и OpenMP, только GLSL и спасает.

[svetlov]

Я бы не рекомендовал. Понимаете, технически это сделать можно. Будет выглядеть как хак на хаке — придется строить геометрию и проч только для того, чтобы запустить шейдер. А потом еще и вытаскивать полученный результат неочевидным образом. Очень неудачное решение.
Если у вас есть CUDA — берите ее. Аналогичные технологии для ATI-AMD, OpenCL, DirectCompute — что угодно. Но не стоит пытаться запускать GPU в графическом режиме для решения вычислительных задач.
Кстати, на старых NVidia, на которых CUDA не идет — не очень много получится с шейдеров выжать (сильно эти карточки специфично под графику заточены). А код нужно будет писать заново, там несовместимостей вагон с тележкой.
И прироста скорости по сравнению с теми же CUDA/OpenCL не будет — скорее наоборот. ALU ведь одни и те же, а выкрутасы с графикой дадут о себе знать.

[Robotex]

А кто сказал, что это под старые карточки Нвидиа? Это под современные карточки Intel. Спасибо, что это объяснили, но не я это придумал :(

[svetlov]

На старых интелах ничего не спасет. Они держат слишком древний GLSL. И имеют проблемы с render target. Рисовать ведь нужно не на экран, а в буфер — чтобы его потом читать. А вот буфера как раз и нет — слишком тупое железо.

[Robotex]

А современные Intel?

[svetlov]

А на новых, насколько помню, OpenCL хромает — но идет. Вопрос небольшого времени, как мне кажется. Летом-осенью станет еще лучше.

[Robotex]

Ну вот, а нужно это именно сейчас. Мы пытались запустить OpenCL — не вышло.

[svetlov]

software.intel.com/en-us/articles/intel-opencl-sdk/

[svetlov]

Сам не пробовал — только OpenCL на NVidia, не больше

[Robotex]

Это только под Windows — не подходит.

[svetlov]

Дело, конечно, ваше — но я бы не рекомендовал. Тем более если, как вы пишите, с OpenGL совсем не знакомы. Долго, мучительно — и результат не радует.

[Robotex]

Боже, да мне всего лишь массивы нужно перемножить.

Answer 5

[svetlov]

Ну смотрите.
Там нет массивов.
Создаете окошко — это обязательно.
Инииализируете конвейер.
Создаете vertex buffer object с координатами вершин. Да, я знаю что вершин у вас нет — но иначе шейдер не запустить, он будет работать по разу на вершину. Вершины могут быть фейковые.
Создаете текстуры. Современные видеокарты поддерживают текстуры из float — проблем быть не должно. Но следует убедиться для той видюхи, на которой все запускается — иначе нужно использовать другой способ расчета.
Устанавливаем все это добро как render state. Запускаем созданный и скомпилированный шейдер.
Он берет значения из текстур и рисует нам картинку. Вернее, шейдеров нужно два: вершинный и пиксельный. Второй должен записать результат. Есть еще выкрутасы с записью промежуточных результатов из шейдера — но тут нужно смотреть, что умеет конкретная видюха.
Берем результат. Откуда — вопрос. Если поддерживаются pixel buffer object и render target — то из PBO. Иначе из буфера экрана (надеюсь, до этого дело не дойдет, видюха будет более или менее сносной).
Может, я еще что-то забыл.

Перечисленная схема — лишь краткое изложение. На самом деле кода получается — мама не горюй. На «простое перемножение массивов» это ни разу не похоже.

Comments

[Robotex]

Мда: pastebin.com/ggETxnX8 выглядит страшно

Answer 6

[Robotex]

CUDA: pastebin.com/TkrhuEWA
Что-то не так. Если L равно 1, то считает верно, если 2 и выше, то какие-то запредельные числа выходят.

for(int j=0;j<L;j++)
        {
          mulKernel<<<blocksMul, threadsMul>>>(devTemp, devInputs, devWeights, j*N*I, I);
          
          sumKernel<<<blocksSum, threadsSum>>>(devInputs, devTemp, N);
        }

Answer 7

[svetlov]

Что-то мне кажется, что без __syncthreads() дела не будет.

Comments

[Robotex]

В каком месте поставить? Между кернелами?

Answer 8

[Robotex]

Здесь: pastebin.com/ggETxnX8 реализовано умножение двух массивов.

У меня несколько вопросов:
1. Количество элементов тут должно быть кратно 4. Как исправить, чтобы работать с любым количеством.
2. Откуда кернел знает с каким индексом он работает?
3. Как использовать два кернела и запускать один после другого?