同一个流内顺序执行以下命令：（1）主机到内存设备拷贝；（2）内核函数；（3）设备到主机内存拷贝。分别在不同命令执行前和执行后压入一个事件记录时间。
我自己在GPU上用cuda平台执行的时候，3个命令的执行时间加起来刚好等于总的运行时间。
在DCU上，就出现了问题，3个命令执行时间加起来不是总的运行时间。不知道是什么原因。
代码如下：

#include <iostream>
#include "hip/hip_runtime.h"
#include <hip/hip_runtime.h>

void initialData(float * v, int size)
{
for(int i = 0; i < size; i++)
v[i] = i;
}

global void add(float *d_A, float *d_B, float *d_C, int N)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
{
d_C[tid] = d_A[tid] * sin(d_B[tid]);
}
}

int main(int argc, char **argv)
{
int ndcus;
hipGetDeviceCount(&ndcus);
printf("HIP-capable devices: %i\n", ndcus);

float *d_A[ndcus], *d_B[ndcus], *d_C[ndcus]; //设备内存

float *h_A[ndcus], *h_B[ndcus], *dcuRef[ndcus];	 //主机内存

hipStream_t stream[ndcus];

//平均分配到各个设备上
int size = 1 << 26;
int iSize = size / ndcus;

std::cout << "个数" << iSize << std::endl;

size_t iBytes = iSize * sizeof(float);

std::cout << "内存大小" << iBytes << std::endl;

//分配主机和设备内存
for (int i = 0; i < ndcus; i++)
{
	// 设置DCU
	hipSetDevice(i);

	// 分配设备内存
	hipMalloc((void **) &d_A[i], iBytes);
	hipMalloc((void **) &d_B[i], iBytes);
	hipMalloc((void **) &d_C[i], iBytes);

	//分配锁页主机内存
	hipHostMalloc((void **) &h_A[i], iBytes);
	hipHostMalloc((void **) &h_B[i], iBytes);
	hipHostMalloc((void **) &dcuRef[i], iBytes);

	// 创建流对象
	hipStreamCreate(&stream[i]);
}

//数据初始化
for (int i = 0; i < ndcus; i++) {
	hipSetDevice(i);
	initialData(h_A[i], iSize);
	initialData(h_B[i], iSize);
}

//for(int i =0; i < ndcus; i++)
//for(int j= 0; j < 5; j++)
// std::cout << h_A[i][j]<< " " << h_B[i][j]<< std::endl;

hipEvent_t memcpy_to_start_event[ndcus];
hipEvent_t kernel_start_event[ndcus];
hipEvent_t memcpy_from_start_event[ndcus]; 
hipEvent_t memcpy_from_stop_event[ndcus]; 

dim3 block(256, 1);
dim3 grid(iSize / 256 + 1, 1);

float time_copy_to_ms = 0.0f;  
float time_kernel_ms = 0.0f; 
float time_copy_from_ms = 0.0f;  
float time_exec_ms = 0.0f;  
  

for(int i = 0; i < ndcus; i++) 
{
	hipSetDevice(i);
  
 hipEventCreate(&memcpy_to_start_event[i]);
 hipEventCreate(&kernel_start_event[i]);
 hipEventCreate(& memcpy_from_start_event[i]);
 hipEventCreate(& memcpy_from_stop_event[i]);
 
  
hipEventRecord(memcpy_to_start_event[i], stream[i]);

//主机拷入设备1
hipMemcpyAsync(d_A[i], h_A[i], iBytes, hipMemcpyHostToDevice, stream[i]);
hipMemcpyAsync(d_B[i], h_B[i],  iBytes, hipMemcpyHostToDevice, stream[i]);

hipEventRecord(kernel_start_event[i], stream[i]);

 
//内核函数
hipLaunchKernelGGL(add, grid, block, 0, stream[i], d_A[i], d_B[i], d_C[i], iSize); 

hipEventRecord(memcpy_from_start_event[i], stream[i]);

//设备考入主机
hipMemcpyAsync(dcuRef[i], d_C[i], iBytes, hipMemcpyDeviceToHost, stream[i]);  

hipEventRecord(memcpy_from_stop_event[i], stream[i]); 
     
}

for(int i = 0; i < ndcus; i++) 
{
   //hipSetDevice(i);
   hipStreamSynchronize(stream[i]);
   
      //内存从设备到主机的拷贝时间

   hipEventElapsedTime(&time_copy_to_ms, memcpy_to_start_event[i], kernel_start_event[i]);
   
   //内存从主机到设备的拷贝时间
   hipEventElapsedTime(&time_kernel_ms, kernel_start_event[i], memcpy_from_start_event[i]);
   
   //内存从设备到主机的拷贝时间
   hipEventElapsedTime(&time_copy_from_ms, memcpy_from_start_event[i],  memcpy_from_stop_event[i]);
   
    //总执行时间
   hipEventElapsedTime(& time_exec_ms, memcpy_to_start_event[i], memcpy_from_stop_event[i]);
 
 float gpu_time = time_copy_to_ms + time_kernel_ms + time_copy_from_ms  ;
 
   std::cout << "device" << i << std::endl;
   std::cout << "Copy to: " << time_copy_to_ms << std::endl;
   std::cout << "Kernel: " << time_kernel_ms << std::endl;
   std::cout << "Copy back: " << time_copy_from_ms << std::endl;
   std::cout << "Component time: " << gpu_time << std::endl;
   std::cout << "Execution time: " << time_exec_ms << std::endl;

   
  std::cout <<  std::endl;   
 }

//等所有设备完毕
hipDeviceSynchronize();

//测试结果
for(int i = 0; i < ndcus; i++)
{
	for(int j = 0; j < 10; j++)
	   std::cout << dcuRef[i][j] << std::endl;


	std::cout << std::endl;
}

//释放内存
std::cout << "释放内存";
for(int i = 0; i < ndcus; i++)
{
	hipFree(d_A[i]);
	hipFree(d_B[i]);
	hipFree(d_C[i]);

	hipHostFree(h_A[i]);
	hipHostFree(h_B[i]);
	hipHostFree(dcuRef[i]);
  
    hipEventDestroy( memcpy_to_start_event[i]);
    hipEventDestroy( kernel_start_event[i]);
     hipEventDestroy( memcpy_from_start_event[i]);
     hipEventDestroy(memcpy_from_stop_event[i]);

}


return 0;

}