Сроки в ArrayFire

Я пытаюсь оценить производительность простых поэлементно-матричных операций графического процессора с помощью ArrayFire.

В частности, учитывая

int N1 = something;
int N2 = something;

array A_D = constant(1.,N1*N2,1,f64);
array B_D = constant(1.,N1*N2,1,f64);
array C_D = constant(1.,N1*N2,1,f64);
array D_D = constant(1.,N1*N2,1,f64);

Я хотел бы выполнить сроки следующей инструкции

D_D = A_D + B_D + C_D + 3.;

Я использую два подхода. Первый

timer  time_last;
time_last = timer::start();

D_D = A_D + B_D + C_D + 3.;

double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);

Второй определяет следующую функцию

void timing_test()
{
    int N1 = something;
int N2 = something;

    array A_D = constant(1.,N1*N2,1,f64);
    array B_D = constant(1.,N1*N2,1,f64);
    array C_D = constant(1.,N1*N2,1,f64);
    array D_D = constant(1.,N1*N2,1,f64);

    D_D = A_D + B_D + C_D + 3.;
}

а потом звонит

printf("elapsed time using timeit %g ms \n", 1000.*timeit(timing_test));

Я получил следующие результаты:

(N1,N2)=(256,256) первый подход = 0.0456ms второй подход = 0.264ms

(N1,N2)=(512,512) первый подход = 0.0451ms второй подход = 0.264ms

(N1,N2)=(1024,1024) первый подход = 0.0457ms второй подход = 0.263ms

(N1,N2)=(2048,2048) первый подход = 0.127ms второй подход = 0.265ms

Я также использую следующую "кодированную" версию выражения в соответствии с

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

eval_matrix_wrap_handcoded(A_D,B_D,C_D,D_D,N1*N2);

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);

template <class T1, class T2, class T3, class T4>
__global__ inline void evaluation_matrix_handcoded(T1 *A_D, T2 *B_D, T3 *C_D, T4 *D_D, int NumElements)
{
    const int i = blockDim.x * blockIdx.x + threadIdx.x;
    if(i < NumElements) D_D[i]=A_D[i]+B_D[i]+C_D[i]+3.;
}

__host__ void eval_matrix_wrap_handcoded(double *A_D, double *B_D, double *C_D, double *D_D, int NumElements)
{
    dim3 dimGrid(iDivUp(NumElements,dimBlock.x));
    evaluation_matrix_handcoded<<<dimGrid,dimBlock>>>(A_D,B_D,C_D,D_D,NumElements);
}

получить следующее

(N1,N2)=(256,256)0.0897ms

(N1,N2)=(512,512)0.339ms

(N1,N2)=(1024,1024)1.3ms

(N1,N2)=(2048,2048)5.37ms

Странно то, что

  1. Результаты двух подходов различны. Это может быть связано с накладными расходами при вызове функции, но в любом случае странно, что эти накладные расходы изменяются, когда (N1,N2)=(2048,2048),
  2. Результаты двух подходов практически не зависят от размеров матрицы.
  3. Результаты сильно отличаются по сравнению с "ручной" версией выражения (я предполагаю, что библиотека должна иметь компромисс между производительностью и производительностью).

Обратите внимание, что перед любой операцией я разогреваю графический процессор с помощью кода

array test1(1,5);
test1(0,0)=1;
test1(0,1)=2;
test1(0,2)=3;
test1(0,3)=4;
test1(0,4)=5;

Может ли кто-нибудь помочь мне интерпретировать вышеупомянутые результаты? Благодарю.

ИЗМЕНИТЬ СЛЕДУЮЩИЙ ОТВЕТ ПАВАНА

Первый метод изменен на

timer  time_last;
time_last = timer::start();

D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();

double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);

Второй метод изменен на

void timing_test()
{
    int N1 = something;
    int N2 = something;

    array A_D = constant(1.,N1*N2,1,f64);
    array B_D = constant(1.,N1*N2,1,f64);
    array C_D = constant(1.,N1*N2,1,f64);
    array D_D = constant(1.,N1*N2,1,f64);

    D_D = A_D + B_D + C_D + 3.;
    D_D.eval();
}

Однако время сейчас

`(N1,N2)=(256,256)`  first approach = `14.7ms`  second approach = `2.04ms`

`(N1,N2)=(512,512)`  first approach = `14.3ms`  second approach = `2.04ms`

`(N1,N2)=(1024,1024)`  first approach = `14.09ms`  second approach = `2.04ms`

`(N1,N2)=(2048,2048)`  first approach = `16.47ms`  second approach = `2.04ms`

и у меня все еще разные тайминги и они не зависят от размера векторов.

Если я изменю первый метод на

D_D = A_D + B_D + C_D + 3.;
D_D.eval();

timer  time_last;
time_last = timer::start();

D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();

double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);

а именно, я "увеличиваю" стадию прогрева графического процессора, получая, для первого метода,

`(N1,N2)=(256,256)`  `0.19ms`

`(N1,N2)=(512,512)`  `0.42ms`

`(N1,N2)=(1024,1024)`  `1.18ms`

`(N1,N2)=(2048,2048)`  `4.2ms`

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

ВТОРОЕ РЕДАКТИРОВАНИЕ Подводя итог: я учел ответ и комментарий и, для первого подхода, я использую

D_D = A_D + B_D + C_D + 3.;
D_D.eval();

timer  time_last;
af::sync();
time_last = timer::start();

D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();

double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);

Я получаю следующие (новые) результаты:

`(N1,N2)=(256,256)`  `0.18ms`

`(N1,N2)=(512,512)`  `0.30ms`

`(N1,N2)=(1024,1024)`  `0.66ms`

`(N1,N2)=(2048,2048)`  `2.18ms`
Источник

1 ответ

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

Это означает что-то вроде

D_D = A_D + B_D + C_D + 3.;

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

Вы можете принудительно оценить эти выражения, если вы используете af::eval() функция или eval() метод.

Так что для вашей конкретной проблемы, пожалуйста, используйте D_D.eval() для обоих методов. Вам нужно будет сделать af::sync() также для первого метода. timeit() не нужно синхронизироваться явно.

Источник
Другие вопросы по тегам