2

私は9600MGTGPUを搭載したMacBookProで小さなC++/スラストプログラム（下記）を実行しており、関数hで時間が費やされている場所を理解することに興味があります。これは、このコードをできるだけ速く実行して、より大きくすることが目標だからです。 NEPSの値。

そのために、私は関数にclock（）呼び出しを散らかしました。

印刷された時間は、ほとんどすべての時間がthrust::reduceに費やされていることを示しています。実際、報告されたthrust :: reduceの時間は、要素ごとにコサインへの3回の呼び出しを呼び出すthrust::transformの時間の数百倍です。なんで？

当然、私は測定された時間に疑いを持っています。報告された時間が類似しているかどうかを確認するために、2回目のthrust::reduceの呼び出しを挿入しました。そうではありません。2回目の呼び出しで報告された時間は、変動がはるかに大きく、小さくなっています。さらに混乱：なぜですか？

また、2つのカーネル呼び出しの代わりに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 usecs\n", 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 = %g\n", 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 = %d\n", CNV_n);
  while (n--) {
    float sum = h(x,t);
    x=x+EU_dt*sum;
    t=t+EU_dt;
  }
  printf("%f\n",x);
}
4

2 に答える 2

1

パフォーマンスのためにアルゴリズムを最適化する場合は、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 呼び出しは必要ありません)。

于 2012-10-21T19:59:30.563 に答える