Запуски ядра CUDA являются асинхронными. Это означает, что элемент управления возвращается потоку хоста, чтобы он мог перейти к следующей строке кода после запуска ядра до того, как ядро даже начало выполнять .
Это покрыто многочисленными вопросами здесь по тегу cuda. Это распространенная ошибка при синхронизации кода CUDA. Это может повлиять на то, как вы проверяете время на код, а также на то, как вы используете обычный код CUDA. Обычным решением является вставка вызова cudaDeviceSynchronize () перед закрытием области синхронизации . Это гарантирует, что все действия CUDA завершены, когда вы закончите измерение времени.
Когда я превратил то, что у вас есть, в законченный код с надлежащими методами синхронизации, код тяги был на самом деле быстрее. Ваш дизайн ядра неэффективен. Вот моя версия вашего кода, работающего на CUDA 10 на Tesla P100, показывающая, что время между двумя случаями примерно одинаково:
$ cat t469.cu
#include <thrust/transform.h>
#include <thrust/complex.h>
#include <thrust/device_ptr.h>
#include <thrust/execution_policy.h>
#include <cuComplex.h>
#include <iostream>
#include <time.h>
#include <sys/time.h>
#include <cstdlib>
#define USECPSEC 1000000ULL
long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
typedef thrust::complex<float> comThr;
struct Complex_Mul_Complex :public thrust::binary_function<comThr, comThr, comThr>
{
__host__ __device__
comThr operator() (comThr a, comThr b) const{
return a*b;
}
};
double cpuSecond(){
long long dt = dtime_usec(0);
return dt/(double)USECPSEC;
}
__global__ void HardamarProductOnDeviceCC(cuComplex *Result, cuComplex *A, cuComplex *B, int N)
{
unsigned int index = threadIdx.x + blockDim.x * blockIdx.x;
if(index < N)
Result[index] = cuCmulf(A[index],B[index]);
}
int main(){
int N = 720896;
cuComplex *d_Data1, *d_Data2;
cudaMalloc(&d_Data1, N*sizeof(d_Data1[0]));
cudaMalloc(&d_Data2, N*sizeof(d_Data2[0]));
// step 1: type convert (cuComplex->thrust)
comThr *thr_temp1 = reinterpret_cast<comThr*>(d_Data1);
thrust::device_ptr<comThr> thr_d_Data1 = thrust::device_pointer_cast(thr_temp1);
comThr *thr_temp2 = reinterpret_cast<comThr*>(d_Data2);
thrust::device_ptr<comThr> thr_d_Data2 = thrust::device_pointer_cast(thr_temp2);
// step 2: product and timing
Complex_Mul_Complex op_dot;
double iStart = cpuSecond(); // timing class
for(int i=0;i<1000;i++){ // loop 1000 times to get accurate time consumption
thrust::transform(thrust::device,thr_d_Data1,thr_d_Data1+N,
thr_d_Data2,thr_d_Data1,op_dot);
}
cudaDeviceSynchronize();
double iElapse = cpuSecond() - iStart;
std::cout << "thrust time consume: " << iElapse <<std::endl;
int blockSize;
int minGridSize;
int gridSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, HardamarProductOnDeviceCC, 0, 0);
gridSize = (N+blockSize-1)/blockSize;
std::cout << "gridsize: " << gridSize << " blocksize: " << blockSize << std::endl;
dim3 grid(gridSize);
dim3 block(blockSize);
iStart = cpuSecond();
for(int i=0;i<1000;i++){
HardamarProductOnDeviceCC<<<grid,block>>>(d_Data1,d_Data1,d_Data2,N);
}
cudaDeviceSynchronize();
iElapse = cpuSecond() - iStart;
std::cout << "kernel time consume: " << iElapse <<std::endl;
}
$ nvcc -o t469 t469.cu
$ ./t469
thrust time consume: 0.033937
gridsize: 704 blocksize: 1024
kernel time consume: 0.0337021
$
Примечание: для того, чтобы я продемонстрировал правильность своего ответа, для меня важно предоставить полный код. Если вам нужна помощь от других, я предлагаю вам предоставить полный код, а не биты и кусочки, которые нужно собрать, а затем преобразовать в полный код, добавив включения и т. Д. Вы можете делать все, что пожелаете Конечно, но если вам будет проще помогать другим, вы можете обнаружить, что вам легче помочь.