推力应用的性能调整

performance tuning a thrust application

本文关键字:性能 调整 应用      更新时间:2023-10-16

我正在我的macbook pro w/9600M GT gpu上运行一个小型C++/推力程序(如下),我有兴趣了解在函数h中花费的时间,因为目标是为了获得更大的NEPS值而尽快运行此代码。

为此,我在函数中加入了clock()调用。

打印的时间表明,几乎所有的时间都花在了推力上::reduce。事实上,报告的推力::reduce的时间是推力::transform的几百倍,后者对每个元素调用三次余弦调用。为什么?

当然,我对时间的测量持怀疑态度。我插入了第二个推力::减少只是为了看看报告的时间是否相似:事实并非如此。为第二次呼叫报告的时间具有更高的方差并且更小。更多的困惑:为什么?

我还尝试过使用thrust::transform_reduce(注释掉了)来代替两个内核调用,这两个调用期望它运行得更快——相反,它慢了4%。为什么?

建议不胜感激!

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <iostream>
#include <stdio.h>
#include <stdint.h>

float NEPS = 6.0;
__device__ float EPS;
__device__ float SQEPS;
__device__ float CNV_win;
__device__ float CNV_dt;
int CNV_n;
float EU_dt;
__host__ __device__ float f(float x,float t){
return x*cos(t)+x*cos(t/SQEPS)+cos(t/EPS);
}
struct h_functor
{
const float x, t;
h_functor(float _x, float _t) : x(_x),t(_t) {}
__host__ __device__
float operator()(const float & t_f) const {
return f(x,   t-CNV_win+CNV_dt*(t_f+1)   )*CNV_dt;
} 
};

clock_t my_clock() __attribute__ ((noinline));
clock_t my_clock() {
return clock();
}
float h(float x,float t){
float sum;
sum = CNV_dt*(f(x,t-CNV_win/2)+f(x,t+CNV_win/2))/2;
clock_t start = my_clock(), diff1, diff2, diff3, diff4, diff5;
thrust::device_vector<float> t_f(CNV_n-2);
diff1 = my_clock() - start;
/* initialize t_f to 0.. CNV_n-3 */
start = my_clock();
thrust::sequence(t_f.begin(), t_f.end());
diff2 = my_clock() - start;
start = my_clock();
thrust::transform(t_f.begin(), t_f.end(), t_f.begin(), h_functor(x,t));
diff3 = my_clock() - start;
start = my_clock();
sum += thrust::reduce(t_f.begin(), t_f.end());
diff4 = my_clock() - start;
start = my_clock();
sum += thrust::reduce(t_f.begin(), t_f.end());
diff5 = my_clock() - start;
#define usec(d) (d)
fprintf(stderr, "Time taken %ld %ld %ld %ld %ld usecsn", usec(diff1), usec(diff2), usec(diff3), usec(diff4), usec(diff5));
/* a bit slower, surprisingly:
sum += thrust::transform_reduce(t_f.begin(), t_f.end(), h_functor(x,t), 0, thrust::plus<float>());
*/
return sum;
}
main(int argc, char ** argv) {
if (argc >= 1) NEPS = strtod(argv[1], 0);
fprintf(stderr, "NEPS = %gn", NEPS);
EPS= powf(10.0,-NEPS);
SQEPS= powf(10.0,-NEPS/2.0);
CNV_win= powf(EPS,1.0/4.0);
CNV_dt = EPS;
CNV_n = powf(EPS,-3.0/4.0);
EU_dt = powf(EPS,3.0/4.0);
cudaMemcpyToSymbol(CNV_win, &CNV_win, sizeof(float));
cudaMemcpyToSymbol(CNV_dt, &CNV_dt, sizeof(float));
cudaMemcpyToSymbol(SQEPS, &SQEPS, sizeof(float));
cudaMemcpyToSymbol(EPS, &EPS, sizeof(float));
float x=1.0;
float t = 0.0;
int n = floor(1.0/EU_dt);
fprintf(stderr, "CNV_n = %dn", CNV_n);
while (n--) {
float sum = h(x,t);
x=x+EU_dt*sum;
t=t+EU_dt;
}
printf("%fn",x);
}

如果您想优化算法的性能,可以选择使用arrayfire。我冒昧地为arrayfire重写了您的代码,您可以将其与推力版本进行比较,并选择运行速度更快的版本:

float h(float x,float t){
float sum = CNV_dt * (f(x, t - CNV_win/2) + f(x, t + CNV_win/2)) / 2;
// initialize t_f with a sequence 0..CNV_n-3
af::array t_f(af::seq(0, CNV_n-3));
// transform vector on the GPU
t_f =  t - CNV_win + CNV_dt*(t_f+1); 
t_f = (x*cos(t_f) + x*cos(t_f/SQEPS) + cos(t_f/EPS)) * CNV_dt;
sum += af::sum<float>(t_f); // sum up all elements of the vector
return sum;
}

此外,请注意,不需要将变量显式复制到GPU(即,不需要cudaMemcpyToSymbol调用)

在多核环境中最好不要使用clock()函数。给出错误的答案是很容易的。

最好使用挂钟clock_gettime。在Windows上,我们也有一些高分辨率计时器。

  • Linux中的高分辨率计时器
  • Windows下的高分辨率计时器
    • "示例代码">

使用CUDA时,最好使用CUDA本身提供的计时器。cutil_timer

  • CUDA设备的初始化成本
  • CUDA 5.0的CUDA计时器