2

我正在尝试使用 ArrayFire 评估简单的 GPU 元素矩阵运算的性能。

特别是，考虑到

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. 与表达式的“手动编码”版本相比，结果有很大不同（我假设库应该在生产力和性能之间进行权衡）。

请注意，在任何操作之前，我正在使用代码预热 GPU

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

有人可以帮我解释上述结果吗？谢谢。

按照 Pavan 的回答进行编辑

第一种方法修改为

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);

即，对于第一种方法，我“增加”了 GPU 预热阶段，我获得了，

`(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`
4

1 回答 1

1

ArrayFire 使用即时编译器进行元素操作（这包括算术、逻辑、三角函数和其他数学运算）。

这意味着类似

D_D = A_D + B_D + C_D + 3.;

存储为表达式，直到用户或另一个非 jit 函数请求 D_D 的值。

如果您使用af::eval()函数或eval()方法，您可以强制评估这些表达式。

因此，对于您的特定问题，请同时使用D_D.eval() 这两种方法。af::sync()对于第一种方法，您也需要这样做。timeit()不需要显式同步。

于 2013-04-03T19:41:19.080 回答