Initialize hip when single kernel is called

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();

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();
+}