From 343cf7e864e53d1b8b59e3de6c0b6979b3270f7a Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Tue, 15 Mar 2016 21:05:15 -0500
Subject: [PATCH] Added single kernel launch to sample

[ROCm/clr commit: a6441c33340cab283555139189036aba9e75f8c8]
---
 .../hipDispatchLatency/hipDispatchLatency.cpp | 169 ++++++++++--------
 1 file changed, 90 insertions(+), 79 deletions(-)

diff --git a/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
index fbdbced063..a47931958f 100644
--- a/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+++ b/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
@@ -36,85 +36,96 @@ __global__ void One(hipLaunchParm lp, float* Ad){
 
 int main(){
 
-hipError_t err;
-float *A, *Ad;
+	hipError_t err;
+	float *A, *Ad;
 
-A = new float[LEN];
-
-for(int i=0;i<LEN;i++){
-	A[i] = 1.0f;
-}
-
-hipStream_t stream;
-err = hipStreamCreate(&stream);
-check("Creating stream",err);
-
-err = hipMalloc(&Ad, SIZE);
-check("Allocating Ad memory on device", err);
-
-err = hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
-check("Doing memory copy from A to Ad", err);
-
-float mS = 0;
-hipEvent_t start, stop;
-hipEventCreate(&start);
-hipEventCreate(&stop);
-
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
-}
-hipDeviceSynchronize();
-hipEventRecord(stop);
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"NULL Stream Sync dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
-
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
-}
-hipEventRecord(stop);
-hipDeviceSynchronize();
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"NULL Stream Async dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
-
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
-hipDeviceSynchronize();
-}
-hipEventRecord(stop);
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"Stream Sync dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
-}
-hipDeviceSynchronize();
-hipEventRecord(stop);
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"Stream Async dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
-
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
-}
-hipEventRecord(stop);
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"NULL Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
-
-hipEventRecord(start);
-for(int i=0;i<ITER;i++){
-hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
-}
-hipEventRecord(stop);
-hipEventElapsedTime(&mS, start, stop);
-std::cout<<"Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
-hipDeviceSynchronize();
+	A = new float[LEN];
 
+	for(int i=0;i<LEN;i++){
+		A[i] = 1.0f;
+	}
+
+	hipStream_t stream;
+	err = hipStreamCreate(&stream);
+	check("Creating stream",err);
+
+	err = hipMalloc(&Ad, SIZE);
+	check("Allocating Ad memory on device", err);
+
+	err = hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+	check("Doing memory copy from A to Ad", err);
+
+	float mS = 0;
+	hipEvent_t start, stop;
+	hipEventCreate(&start);
+	hipEventCreate(&stop);
+
+	hipEventRecord(start);
+	hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"First Kernel Launch: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	
+	hipEventRecord(start);
+	hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"Second Kernel Launch: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
+	}
+	hipDeviceSynchronize();
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"NULL Stream Sync dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
+
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
+	}
+	hipEventRecord(stop);
+	hipDeviceSynchronize();
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"NULL Stream Async dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
+
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
+		hipDeviceSynchronize();
+	}
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"Stream Sync dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
+	}
+	hipDeviceSynchronize();
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"Stream Async dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
+
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
+	}
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"NULL Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
+
+	hipEventRecord(start);
+	for(int i=0;i<ITER;i++){
+		hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
+	}
+	hipEventRecord(stop);
+	hipEventElapsedTime(&mS, start, stop);
+	std::cout<<"Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
+	hipDeviceSynchronize();
 }