настройка производительности приложения тяги

Я использую небольшую программу C ++ / Thrust (ниже) на моем MacBook Pro с 9600M GT GPU и мне интересно понять, на что тратится время в функции h, потому что цель состоит в том, чтобы как можно быстрее запустить этот код для больших значения NEPS.

Для этой цели я засорял функцию вызовами clock ().

Напечатанные времена указывают, что почти все время потрачено на тягу: уменьшение.
На самом деле, сообщаемое время для thrust :: lower в несколько сотен раз больше, чем для thrust :: transform, который вызывает три вызова косинуса на элемент. Зачем?

Естественно, я с подозрением отношусь к измеренным временам.
Я вставил 2-й вызов в thrust :: Reducer, чтобы посмотреть, будет ли указанное время аналогичным: это не так. Время, сообщаемое для 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);
}