Layout the cooperative group related headers to target Nvidia

Change-Id: I530ffc72d6fae1e273b5322d508629221e11ea6f


[ROCm/hip commit: 4b6d92798f]

projects/hip/include/hip/hip_cooperative_groups.h

@@ -30,12 +30,15 @@ THE SOFTWARE.
 #ifndef  HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
 #define  HIP_INCLUDE_HIP_HIP_COOPERATIVE_GROUP_H
 
+#include <hip/hip_version.h>
+#include <hip/hip_common.h>
+
 #if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
 #if __cplusplus && defined(__clang__) && defined(__HIP__)
 #include <hip/hcc_detail/hip_cooperative_groups.h>
 #endif
 #elif defined(__HIP_PLATFORM_NVCC__) && !defined(__HIP_PLATFORM_HCC__)
-#include <cooperative_groups.h>
+#include <hip/nvcc_detail/hip_cooperative_groups.h>
 #else
 #error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
 #endif

projects/hip/include/hip/nvcc_detail/hip_cooperative_groups.h

@@ -0,0 +1,12 @@
+#ifndef HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H
+#define HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H
+
+// Include CUDA headers
+#include <cuda_runtime.h>
+#include <cooperative_groups.h>
+
+// Include HIP wrapper headers around CUDA
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+
+#endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_COOPERATIVE_GROUPS_H