OpenCL код работает на MBP быстрее, чем на NVIDIA GTX 480 - PullRequest
Новая
       16

OpenCL код работает на MBP быстрее, чем на NVIDIA GTX 480

3 голосов
/ 25 мая 2011

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

__kernel void matrix_mult(__global float* a, 
              __global float* b, 
              __global float* c,
              const int N) 
{
  int row = get_global_id(1);
  int col = get_global_id(0);
  float sum = 0.0f;
  for (int i = 0; i < N; i++) {
    sum += a[row*N+i] * b[i*N+col];
  }
  c[row*N+col] = sum;
}

Я тестирую аппаратное обеспечение, выполняя код 100 раз так: 

  clock_t begin=clock(); 

  const unsigned int repeats = 100;
  for(int  i = 0; i != repeats; i++){
    runCL(a, b, results,N, N*N);
  }

  clock_t end=clock();

На моем MBP matrix_multiplications требуется около 1,2 мс, на матрицах размером 512 * 512, в то время как тот же код занимает около 3 мс при работе на Linux GTX 480. Это беспокоит меня с тех пор, я не ожидал бы, что дорогая карта GTX будет немного быстрее, чем ноутбук.

Насколько я вижу, мой код «неправильный», потому что я неправильно рассчитал время.

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

cl_event event = {0}; 
err = clEnqueueNDRangeKernel(cmd_queue, kernel[0], 2, NULL, global_work_size, NULL, 0, NULL, &event);
assert(err == CL_SUCCESS);


cl_int err =  clWaitForEvents (1,&event);
cl_ulong start, end; 
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END,   sizeof(cl_ulong), &end,   NULL); 
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); 
double executionTimeInMilliseconds = (end - start) * 1.0e-6f;
std::cout << "execution time in milis : " << executionTimeInMilliseconds << std::endl;

Теперь GT330M будет выполнять операцию за 46 мс, а GTX480 - за 2,5 мс. Тогда возникает еще один действительно интересный вопрос: при включенном PROFILING GT 330M становится примерно в 30 раз медленнее, это имеет смысл, но GTX480 сохраняет ту же производительность. Кто-нибудь может объяснить, почему это так?

Ответы [ 3 ]

4 голосов
/ 25 мая 2011

С точки зрения исходной проблемы, здесь вы видите, что с этим наивным кодом лучшие характеристики GTX480 действительно причиняют вам боль.

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

GTX480 имеет в 3 раза большую (384 бит) и 2x более быструю (1840 МГц) шину памяти, чем GT330M (128 бит, 800 МГц).Номинально, это дает пиковое преимущество в полосе пропускания 177,4 ГБ / с против 25,6 ГБ / с, и, поскольку здесь преобладает пропускная способность памяти, вы можете подумать, что это победит.Однако из-за не слитых чтений и более широкой шины памяти доступ к b-массиву использует только 32 бита из этого 384-битного доступа к памяти, а в случае 330M только 32 бита из каждого 128-битного доступа.Таким образом, эффективная пропускная способность памяти для доступа b составляет 14,8 ГБ / с и 6,4 ГБ / с;так что теперь разница в общей пропускной способности памяти в 2 раза больше, чем в 7, и так много преимуществ более быстрой карты тратится впустую;Кроме того, эта пропускная способность памяти должна быть поделена в 10 раз на большее количество ядер, поэтому задержка для каждого ядра, чтобы получить доступ и выполнить вычисления, больше.Я подозреваю, что если бы вы использовали матрицы большего размера, вы могли бы скрыть большую часть задержки и приблизиться к максимально возможному 2-кратному ускорению, а не к 2,5-кратному замедлению, которое вы видите.

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

Результаты профилирования, которые вы видите, однако я понятия не имею.Возможно, 330M не имеет такой хорошей аппаратной поддержки для профилирования, так что все должно быть реализовано в программном обеспечении?Так как числа GTX в любом случае примерно одинаковы, я бы сейчас использовал более простой подход синхронизации, который, поскольку вы не используете асинхронные ядра или передачу, должен быть в порядке.

2 голосов
/ 25 мая 2011

Я думаю, что вы раздвигаете пределы разрешения таймера для Nvidia. Попробуйте clGetDeviceInfo () с CL_DEVICE_PROFILING_TIMER_RESOLUTION, чтобы проверить это. С этими крошечными временами я бы ничего не закончил.

1 голос
/ 25 мая 2011

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

...