[hip] add initial implementation for hipLaunchCooperativeKernel API (#1339)

* [hip] add initial implementation for hipLaunchCooperativeKernel API

* [hip] use total number of work groups to initialize the GWS resource

* [hip] use only one argument for init_gws kernel

* [hip] use the device associated with the stream for checking the device properties


[ROCm/hip commit: 5066700ace]

projects/hip/include/hip/hcc_detail/device_library_decls.h

@@ -59,6 +59,10 @@ extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float);
 extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
 extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
 
+extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid);
+extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid);
+
+
 // Introduce local address space
 #define __local __attribute__((address_space(3)))