Improve HIP kernel names, attributes and codegen, contributed by Alex Voicu

Change-Id: I2cafbdc5a98e26c7f4fad84739c915e7dc09993c


[ROCm/clr commit: 3b6a863eef]

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -149,8 +149,15 @@ extern int HIP_TRACE_API;
 #endif /* Device feature flags */
 
 
-//TODO-HCC  this is currently ignored by HCC target of HIP
-#define __launch_bounds__(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)
+#define launch_bounds_impl0(requiredMaxThreadsPerBlock)\
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
+#define launch_bounds_impl1(\
+    requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)\
+    __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock),\
+        amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
+#define select_impl_(_1, _2, impl_, ...) impl_
+#define __launch_bounds__(...) select_impl_(\
+    __VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0)(__VA_ARGS__)
 
 // Detect if we are compiling C++ mode or C mode
 #if defined(__cplusplus)

projects/clr/hipamd/include/hip/hcc_detail/host_defines.h

@@ -48,7 +48,7 @@ THE SOFTWARE.
 #define __global__  __attribute__((hc_grid_launch)) __attribute__((used))
 #else
 //#warning "GGL global define reached"
-#define __global__ __attribute__((hc, weak))
+#define __global__ __attribute__((annotate("hip__global__"), hc, used))
 #endif //GENERIC_GRID_LAUNCH
 
 #define __noinline__      __attribute__((noinline))