Умножение матрицы OpenCL должно быть быстрее?

Question

Я пытаюсь узнать, как сделать GPU оптимальным ядерным ядрам OpenCL, я взял пример умножения матрицы, используя квадратные плитки в локальной памяти. Однако Я получил в лучшем случае только ~ 10-кратное ускорение (~ 50 Gflops) по сравнению с numpy.dot() (5 Gflops, он использует BLAS).

Я нашел исследования, где получили ускорение> 200x (> 1000 Gflops). ftp://ftp.u-aizu.ac.jp/u-aizu/doc/Tech-Report/2012/2012-002.pdf Я не знаю, что я делаю неправильно, или просто из-за моего GPU (nvidia GTX 275). Или это связано с некоторыми издержками pyOpenCl. Но я также оценил, сколько времени требуется для копирования результата с GPU в RAM, и это всего лишь ~ 10% от времени умножения матрицы.

#define BLOCK_SIZE 22 
__kernel void matrixMul(
     __global float* Cij, 
     __global float* Aik, 
     __global float* Bkj, 
     __const int ni, 
     __const int nj, 
     __const int nk 
){ 
// WARRNING : interchange of i and j dimension lower the performance >2x on my nV GT275 GPU  
int gj = get_global_id(0); int gi = get_global_id(1); 
int bj = get_group_id(0);  int bi = get_group_id(1); // Block index 
int tj = get_local_id(0);  int ti = get_local_id(1); // Thread index 
int oj = bi*BLOCK_SIZE;  int oi = bj*BLOCK_SIZE; 
float Csub =0; 
__local float As [BLOCK_SIZE][BLOCK_SIZE]; 
__local float Bs [BLOCK_SIZE][BLOCK_SIZE]; 
for (int ok = 0; ok < nk; ok += BLOCK_SIZE) { 
    As[ti][tj] = Aik[ nk*(gi ) + tj + ok ]; // A[i][k] 
    Bs[ti][tj] = Bkj[ nj*(ti+ok) + gj ];  // B[k][j] 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for (int k = 0; k < BLOCK_SIZE; ++k) Csub += As[ti][k] * Bs[k][tj]; 
    barrier(CLK_LOCAL_MEM_FENCE); 
} 
Cij[ nj * (gi) + gj ] = Csub; 

}

Примечание - странный BLOCK_SIZE = 22 является максимальным BLOCK_SIZE, который делает подходит к максимальному work_group_size, который 512 на моем GPU. В этом коде должно быть указано условие BLOCK_SIZE^2 < max work_group_size. 22 = INT (SQRT (512)). Я попробовал также BLOCK_SIZE = 16 или 8, но был медленнее tan с 22.

Я также пробовал простую матрицуMul (без использования локальной памяти), но это было даже в 10 раз медленнее, чем numpy.dot(). Я скопировал здесь код http://gpgpu-computing4.blogspot.cz/2009/10/matrix-multiplication-3-opencl.html говорят, что даже простая версия (без локальной памяти) должна работать на 200 раз быстрее, чем у процессора? Я этого не делаю.

dependecne производительности в моем случае:

N = 220 numpy 3.680 [Gflops] GPU 16.428 [Gflops] speedUp 4.464 
N = 330 numpy 4.752 [Gflops] GPU 29.487 [Gflops] speedUp 6.205 
N = 440 numpy 4.914 [Gflops] GPU 37.096 [Gflops] speedUp 7.548 
N = 550 numpy 3.849 [Gflops] GPU 47.019 [Gflops] speedUp 12.217 
N = 660 numpy 5.251 [Gflops] GPU 49.999 [Gflops] speedUp 9.522 
N = 770 numpy 4.565 [Gflops] GPU 48.567 [Gflops] speedUp 10.638 
N = 880 numpy 5.452 [Gflops] GPU 44.444 [Gflops] speedUp 8.152 
N = 990 numpy 4.976 [Gflops] GPU 42.187 [Gflops] speedUp 8.478 
N = 1100 numpy 5.324 [Gflops] GPU 83.187 [Gflops] speedUp 15.625 
N = 1210 numpy 5.401 [Gflops] GPU 57.147 [Gflops] speedUp 10.581 
N = 1320 numpy 5.450 [Gflops] GPU 48.936 [Gflops] speedUp 8.979 

Примечание - номер «Gflops» получается как N^3/время, и это не включает в себя время, необходимое для копирования результатов от GPU к оперативной памяти, но на этот раз всего лишь несколько процентов от общего времени, особенно для N> 1000

может быть более изобразительным время в secons:

N = 220 numpy 0.003 [s] GPU 0.001 [s] load 0.001 [s] speedUp 5.000 
N = 330 numpy 0.008 [s] GPU 0.001 [s] load 0.001 [s] speedUp 7.683 
N = 440 numpy 0.017 [s] GPU 0.002 [s] load 0.001 [s] speedUp 7.565 
N = 550 numpy 0.043 [s] GPU 0.004 [s] load 0.001 [s] speedUp 11.957 
N = 660 numpy 0.055 [s] GPU 0.006 [s] load 0.002 [s] speedUp 9.298 
N = 770 numpy 0.100 [s] GPU 0.009 [s] load 0.003 [s] speedUp 10.638 
N = 880 numpy 0.125 [s] GPU 0.010 [s] load 0.000 [s] speedUp 12.097 
N = 990 numpy 0.195 [s] GPU 0.015 [s] load 0.000 [s] speedUp 12.581 
N = 1100 numpy 0.250 [s] GPU 0.031 [s] load 0.000 [s] speedUp 8.065 
N = 1210 numpy 0.328 [s] GPU 0.031 [s] load 0.000 [s] speedUp 10.581 
N = 1320 numpy 0.422 [s] GPU 0.047 [s] load 0.000 [s] speedUp 8.979 

Я думал, что, возможно, некоторые скорости импровизации ement можно получить, используя async_work_group_copy или даже read_imageui для копирования блоков в локальную память. Но я не понимаю, почему у меня такая большая разница, когда я использую в основном тот же код, что и люди, которые говорят, что у них 200-кратное ускорение ?????

Answer 1

Не смотря на свой код, позвольте мне высказать некоторые замечания по поводу ваших тестов. Давайте проигнорируем numpy и сравним максимальные SP FLOPs/s и DP FLOPs/s процессора Intel против графических процессоров Nvidia и AMD.

A Intel 2600K на частоте 4 ГГц может выполнять 4 ГГц * (8 AVX) * (2 ILP) * (4 ядра) = 256 SP GFLOPs/s. Для DP это половина: 128 DP GFLOPs/s. Хасуэлл, который выйдет через несколько недель, удвоит их обоих. Библиотека Intel MKL получает более 80% эффективности в GEMM. Мой собственный код GEMM получает 70% на моем i7-2700, поэтому 5 GFlops/s, которые вы цитируете с numpy, являются крошечными и нечестными для сравнения.

Я не знаю, на что способен GTX 275, но я бы предположил, что это намного больше, чем 50 GFLOP/s.

В данной статье сравнивается AMD 7970. Они получают 848 (эффективность 90%) DP GFlops/s и 2646 (70% эффективности) SP GFlops/s. Это ближе к производительности в 10 раз от CPU не 200x!

Редактировать: Ваши расчеты FLOPs неверны, это должно быть 2.0 * n^3. Это все еще приближенно, но это асимптотически верно. Позволь мне объяснить.

Рассмотрите 3D-продукт. Это x1 * x2 + y1 * y2 + z1 * z2. Это 3 умножения и два дополнения. Таким образом, N-мерный точечный продукт представляет собой n умножений и (n-1) дополнений. Матричный продукт эквивалентен nxn-точечным произведениям, т. Е. N * n * n умножений и n * n * (n-1) дополнений. Это примерно 2.0 * n^3 FLOPS. Таким образом, вы должны удвоить все свои числа Gflops/s.

Редактировать: Возможно, вы захотите рассмотреть время ядра. Прошло некоторое время, поскольку я использовал OpenCL, но с использованием C++ привязок я сделал что-то вроде этого

queue = cl::CommandQueue(context, devices[device], CL_QUEUE_PROFILING_ENABLE|CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err); 
//other code...run kernel 

time_end = clevent.getProfilingInfo<CL_PROFILING_COMMAND_END>(); 
time_start = clevent.getProfilingInfo<CL_PROFILING_COMMAND_START>(); 

Answer 2

Хороший GPU матрицы умножения не просто использовать локальную память, он хранит блоки А, В и/или С в регистры (что приводит к более высокому использованию регистров и снижению занятости, но в конце концов намного быстрее). Это связано с тем, что у графических процессоров больше регистров, чем у локальной памяти (128-256 КБ против 48 КБ для NVIDIA), а регистры предлагают такую ​​же пропускную способность, что и ALU.