From 5efc4ce08a989db4e247429dcc6d5b635ad018d1 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Wed, 2 Mar 2016 08:08:45 -0600
Subject: [PATCH] Initialize hip when single kernel is called

---
 hipamd/src/hip_hcc.cpp         | 1 +
 hipamd/tests/src/hipKernel.cpp | 8 ++++++++
 2 files changed, 9 insertions(+)
 create mode 100644 hipamd/tests/src/hipKernel.cpp

diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp
index 09b4770347..1cf8745c0b 100644
--- a/hipamd/src/hip_hcc.cpp
+++ b/hipamd/src/hip_hcc.cpp
@@ -1090,6 +1090,7 @@ inline hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 // Allows runtime to track some information about the stream.
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av)
 {
+	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
 
     stream->preKernelCommand();
diff --git a/hipamd/tests/src/hipKernel.cpp b/hipamd/tests/src/hipKernel.cpp
new file mode 100644
index 0000000000..6bf089bc4e
--- /dev/null
+++ b/hipamd/tests/src/hipKernel.cpp
@@ -0,0 +1,8 @@
+#include"test_common.h"
+
+__global__ void Empty(hipLaunchParm lp, int param){}
+
+int main(){
+hipLaunchKernel(HIP_KERNEL_NAME(Empty), dim3(1), dim3(1), 0, 0, 0);
+hipDeviceSynchronize();
+}