Оптимизация OpenACC против GPU

Я нуб в OpenACC, я пытаюсь оптимизировать код, для процессора я получаю:

Time = Time + omp_get_wtime();
{
#pragma acc parallel loop
for (int i = 1;i < k-1; i++)
{
jcount[i]=((int)(MLT[i]/dt))+1;
}
jcount[0]=0;
jcount[k-1]=N;

#pragma acc parallel loop collapse(2)
for (int i = 0;i < k - 1; i++)
{
for(int j=jcount[i];j < jcount[i+1];j++)
{
w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
}
}
}
Time = omp_get_wtime() - Time;

Для моего Intel I7 (я отключил Hyper-Threading) с 6 ядрами, я получаю плохое распараллеливание, разница между ядрами 6 против 1 только 30% (это означает, что 70% кода выполняется последовательно, но я не вижу, где)

Для графического процессора:

...
acc_init( acc_device_nvidia );
...
TimeGPU = TimeGPU + omp_get_wtime();
{
#pragma acc kernels loop independent  copyout(jcount[0:k]) copyin(MLT[0:k],dt)
for (int i = 1;i < k-1; i++)
{
jcount[i]=((int)(MLT[i]/dt))+1;
}
jcount[0]=0;
jcount[k-1]=N;

#pragma acc kernels loop independent copyout(X[0:N+1],Y[0:N+1]) copyin(MLT[0:k],MLX[0:k],MLY[0:k],dt) copy(w[0:N])
for (int i = 0;i < k - 1; i++)
{
for(int j=jcount[i];j < jcount[i+1];j++)
{
w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
}
}
}
TimeGPU = omp_get_wtime() - TimeGPU;

А также ГПУ (gtx1070) 3 раза медленнее чем 6 ядер процессора!

Launch parameters:
GPU: pgc++ -ta=tesla:cuda9.0 -Minfo=accel -O4
CPU: pgc++ -ta=multicore -Minfo=accel -O4

k = 20000, N = 2 миллиона

ОБНОВИТЬ:

изменить код графического процессора:

TimeGPU = TimeGPU + omp_get_wtime();
#pragma acc data create(jcount[0:k],w[0:N]) copyout(X[0:N+1],Y[0:N+1]) copyin(MLT[0:k],MLX[0:k],MLY[0:k],dt)
{
#pragma acc parallel loop
for (int i = 1;i < k-1; i++)
{
jcount[i]=((int)(MLT[i]/dt))+1;
}
jcount[0]=0;
jcount[k-1]=N;

#pragma acc parallel loop
for (int i = 0;i < k - 1; i++)
{
for(int j=jcount[i];j < jcount[i+1];j++)
{
w[j] = (j*dt - MLT[i])/(MLT[i+1]-MLT[i]);
X[j] = MLX[i]*(1-w[j])+MLX[i+1]*w[j];
Y[j] = MLY[i]*(1-w[j])+MLY[i+1]*w[j];
}
}
}
TimeGPU = omp_get_wtime() - TimeGPU;
Launch parameters:
pgc++ -ta=tesla:managed:cuda9.0 -Minfo=accel -O4

Теперь GPU в 2 раза медленнее этого процессора

Выход:

139: compute region reached 1 time
139: kernel launched 1 time
grid: [157]  block: [128]
device time(us): total=425 max=425 min=425 avg=425
elapsed time(us): total=509 max=509 min=509 avg=509
139: data region reached 2 times
139: data copyin transfers: 1
device time(us): total=13 max=13 min=13 avg=13
146: compute region reached 1 time
146: kernel launched 1 time
grid: [157]  block: [128]
device time(us): total=13,173 max=13,173 min=13,173 avg=13,173
elapsed time(us): total=13,212 max=13,212 min=13,212 avg=13,212

Почему я получаю TimeGPU в 2 раза больше по сравнению с выводом с использованием PGI_ACC_TIME = 1?
(30 мс против 14 мс)