make hip-tests compileable with TheRock (#1624)

## Motivation Resolved: SWDEV-566226 The current implementation of agents inside of rocprof-systems keeps just the minimal necessary set of information required for populating the `info_agent` table inside of rocpd database. There is a sufficient amount of data that is being left out from database, so this change should fix that and store the additional agent information as an `extdata` row inside of `info_agent` table. ## Technical Details This PR introduces additional filed inside of `agent` structure inside which is representing the JSON formatted string of all the additional information we can acquire about particular agent. This data is processed and added during the initial fetching of agents, and afterwards pushed inside of the database. --------- Co-authored-by: David Galiffi <David.Galiffi@amd.com> * SWDEV-557412 - Incorporate proper chunk offset when remapping virtual memory (#1848) * SWDEV-557412 - Incorporate proper offset when remapping virtual memory * Fix condition to check if VMHeap allocation address matches a chunk address * Move offset calculation outside if/else block --------- Co-authored-by: JeniferC99 <150404595+JeniferC99@users.noreply.github.com> * SWDEV-567852 - Clean-up hip::init() (#1948) * SWDEV-559267 - Use CLPrint to DevLogPrintf with Log Level - detail debug. (#1160) * SWDEV-548892 - Stop using ocml isinf wrapper (#1854) * SWDEV-562708 - change default maximum SVM size to 256GB (#1731) * SWDEV-503089 - Fix and enable disabled HIP tests from math group (#1319) * SWDEV-503089 - Fix and enable disabled HIP tests from math group * SWDEV-503089 - Move single precision reduced run to a common function * SWDEV-548892 - Stop using ockl steadyctr function (#1882) Directly use the builtin * Implement PTL support (#1957) * Implement PTL support Signed-off-by: adapryor <Adam.pryor@amd.com> (cherry picked from commit 45bc31292e7940a3b8fca044ef7df22047b95733) Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> --------- Signed-off-by: adapryor <Adam.pryor@amd.com> Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> Co-authored-by: Maisam Arif <Maisam.Arif@amd.com> * SWDEV-558080 - Add recommended granularity (#1176) * Add recommended granularity * Improve granularity testing * Update based on feedback * Fix and enable VMM tests on cuda (#1855) * Fix and enable VMM tests on cuda * Minor syntax fixes --------- Co-authored-by: Rahul Manocha <rmanocha@amd.com> * [rocprofiler-systems] Add support for ompt_callback_thread_begin (#1681) * Add thread_begin callback * Make OMPT callbacks that are instant have start_ts = end_ts * SWDEV-567514: Remove default stream wait (#1977) - when virtual map command is called - can create deadlock Signed-off-by: sdashmiz <shadi.dashmiz@amd.com> * Fix flaky test Unit_hipStreamAddCallback_StrmSyncTiming (#2022) * Review comments * skip the 3 failing tests to merge hip-tests rocm-systems PR --------- Signed-off-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> Signed-off-by: adapryor <Adam.pryor@amd.com> Signed-off-by: Maisam Arif <Maisam.Arif@amd.com> Signed-off-by: sdashmiz <shadi.dashmiz@amd.com> Co-authored-by: GunaShekar <agunashe@amd.com> Co-authored-by: agunashe <ajay.gunashekar@amd.com> Co-authored-by: Ethan Trinh <Ethan.Trinh@amd.com> Co-authored-by: JeniferC99 <150404595+JeniferC99@users.noreply.github.com> Co-authored-by: Victor Zhang <111778801+victzhan@users.noreply.github.com> Co-authored-by: German Andryeyev <56892148+gandryey@users.noreply.github.com> Co-authored-by: usrihari123 <srihari.u@amd.com> Co-authored-by: Bindhiya Kanangot Balakrishnan <Bindhiya.KanangotBalakrishnan@amd.com> Co-authored-by: anujshuk-amd <anujshuk@amd.com> Co-authored-by: itrowbri <Ian.Trowbridge@amd.com> Co-authored-by: marantic-amd <marantic@amd.com> Co-authored-by: David Galiffi <David.Galiffi@amd.com> Co-authored-by: cadolphe-amd <chris.adolphe@amd.com> Co-authored-by: Karthik Jayaprakash <54370791+kjayapra-amd@users.noreply.github.com> Co-authored-by: Matt Arsenault <Matthew.Arsenault@amd.com> Co-authored-by: Todd tiantuo Li <88386084+lttamd@users.noreply.github.com> Co-authored-by: amilanov-amd <Aleksandar.Milanov@amd.com> Co-authored-by: Adam Pryor <61172547+adam360x@users.noreply.github.com> Co-authored-by: Maisam Arif <Maisam.Arif@amd.com> Co-authored-by: AidanBeltonS <abeltons@amd.com> Co-authored-by: Rahul Manocha <153310294+manocharahul@users.noreply.github.com> Co-authored-by: Rahul Manocha <rmanocha@amd.com> Co-authored-by: Kian Cossettini <Kian.Cossettini@amd.com> Co-authored-by: Shadi Dashmiz <94885391+shadidashmiz@users.noreply.github.com> Co-authored-by: Ioannis Assiouras <38722728+iassiour@users.noreply.github.com> Co-authored-by: Ajay GunaShekar <86270081+agunashe@users.noreply.github.com>
2025-12-03 16:53:17 +00:00
parent f3ffd7070c
commit 8e1aee62d0
124 changed files with 1203 additions and 19124 deletions
@@ -39,39 +39,55 @@ set(TEST_SRC
 )

 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code
-                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc
+                   COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                   -x hip ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc
+                    -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                   -o get_function_module.code
-                   -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc)
 add_custom_target(get_function_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code)
+#target_link_libraries(get_function_module.code hip::host hip::device)
+
+if (NOT WIN32)
+  add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
+                     COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                     -x hip ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc
+                     -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
+                     -o launch_kernel_module.code
+                     DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc)
+  #target_link_libraries(launch_kernel_module.code hip::host hip::device)
+  add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code)
+
+  add_custom_target(coopKernel.code
+                    COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                    -x hip ${CMAKE_CURRENT_SOURCE_DIR}/coopKernel.cpp
+                    -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/coopKernel.code
+                    -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
+                    -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+  #target_link_libraries(coopKernel.code hip::host hip::device)
+  set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS 
+  ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
+  ${CMAKE_CURRENT_BINARY_DIR}/coopKernel.code)
+endif()
+

-add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
-                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc
-                   -o launch_kernel_module.code
-                   -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc)
-add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code)

 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code
-                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc
+                   COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                   -x hip ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc
+                   -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                   -o get_global_test_module.code
-                   -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc)
 add_custom_target(get_global_test_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code)
+#target_link_libraries(get_global_test_module.code hip::host hip::device)

 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code
-                   COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc
+                   COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                   -x hip ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc
+                   -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                   -o get_tex_ref_module.code
-                   -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc)
 add_custom_target(get_tex_ref_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code)
-
-add_custom_target(coopKernel.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/coopKernel.cpp
-                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/coopKernel.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(get_tex_ref_module.code hip::host hip::device)

 add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt
                   COMMAND ${CMAKE_COMMAND} -E copy
@@ -87,18 +103,15 @@ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt
 add_custom_target(empty_file ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt)

 add_custom_target(empty_module
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/empty_module.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})

 set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code
-  ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code
-  ${CMAKE_CURRENT_BINARY_DIR}/coopKernel.code
  ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt
  ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt
  ${CMAKE_CURRENT_BINARY_DIR}/empty_module.code
@@ -120,67 +133,91 @@ if(BUILD_SHARED_LIBS)
 endif()

 add_custom_target(copyKernel.code
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=5 --cuda-device-only
+                  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(copyKernel.code hip::host hip::device)

 add_custom_target(copyKernel.s
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 -S ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=5 -S -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.s
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(copyKernel.s hip::host hip::device)

 add_custom_target(addKernel.code
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=5 --cuda-device-only -x hip ${OFFLOAD_ARCH_LIST}
                  ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/addKernel.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(addKernel.code hip::host hip::device)  

-add_custom_target(addKernel.spv
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco --offload-arch=amdgcnspirv
-                  --no-gpu-bundle-output
-                  ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
-                  -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+# TODO - Rock build failure from clang
+if(NOT WIN32)
+  add_custom_target(addKernel.spv
+                    COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only --offload-arch=amdgcnspirv
+                    --no-gpu-bundle-output
+                    -x hip ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
+                    -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv
+                    -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
+                    -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+  #target_link_libraries(addKernel.spv hip::host hip::device)
+
+  add_custom_target(addKernel-bundle.spv
+                    COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only --offload-arch=amdgcnspirv
+                    -x hip ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
+                    -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv
+                    -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
+                    -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+  #target_link_libraries(addKernel-bundle.spv hip::host hip::device)
+  set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
+    ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv
+    ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv)
+endif()

-add_custom_target(addKernel-bundle.spv
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco --offload-arch=amdgcnspirv
-                  ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc
-                  -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)

 add_custom_target(copyKernelCompressed.code
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --offload-compress --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=5 --offload-compress --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelCompressed.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(copyKernelCompressed.code hip::host hip::device)

-set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic")
+set(OFFLOAD_ARCH_GENERIC_STR
+  --offload-arch=gfx9-generic
+  --offload-arch=gfx9-4-generic:sramecc+:xnack-
+  --offload-arch=gfx9-4-generic:sramecc-:xnack-
+  --offload-arch=gfx9-4-generic:xnack+
+  --offload-arch=gfx10-1-generic
+  --offload-arch=gfx10-3-generic
+  --offload-arch=gfx11-generic
+  --offload-arch=gfx12-generic
+)
+#set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic")
 add_custom_target(copyKernelGenericTarget.code
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --genco ${OFFLOAD_ARCH_GENERIC_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=6 --cuda-device-only ${OFFLOAD_ARCH_GENERIC_STR}
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTarget.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(copyKernelGenericTarget.code hip::host hip::device)

 add_custom_target(copyKernelGenericTargetCompressed.code
-                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --offload-compress --genco ${OFFLOAD_ARCH_GENERIC_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} -mcode-object-version=6 --offload-compress --cuda-device-only
+                  ${OFFLOAD_ARCH_GENERIC_STR}
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTargetCompressed.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${HIP_INCLUDE_DIR} ${HIP_PATH_OPT}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+#target_link_libraries(copyKernelGenericTargetCompressed.code hip::host hip::device)
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.code
  ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.s
  ${CMAKE_CURRENT_BINARY_DIR}/addKernel.code
-  ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv
-  ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv
  ${CMAKE_CURRENT_BINARY_DIR}/copyKernelCompressed.code
  ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTarget.code
  ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTargetCompressed.code
@@ -192,59 +229,60 @@ set(TEST_SRC
    hipKerArgOptimization.cc)

 add_custom_target(copiousArgKernel.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel.code hip::host hip::device)

 add_custom_target(copiousArgKernel0.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=0
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel0.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel0.code hip::host hip::device)

 add_custom_target(copiousArgKernel1.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=1
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel1.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel1.code hip::host hip::device)

 add_custom_target(copiousArgKernel2.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=2
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel2.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel2.code hip::host hip::device)

 add_custom_target(copiousArgKernel3.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=3
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel3.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel3.code hip::host hip::device)

 add_custom_target(copiousArgKernel16.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=16
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel16.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel16.code hip::host hip::device)

 add_custom_target(copiousArgKernel17.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only ${OFFLOAD_ARCH_LIST}
                  -mllvm -amdgpu-kernarg-preload-count=17
-                  ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
+                  -x hip ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc
                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel17.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(copiousArgKernel17.code hip::host hip::device)
+
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel.code
  ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel0.code
@@ -258,7 +296,7 @@ endif()
 endif()

 if(HIP_PLATFORM MATCHES "amd")
-    set(RTCLIB "hiprtc")
+    set(RTCLIB "hiprtc::hiprtc")
 else()
    set(RTCLIB "nvrtc")
 endif()
@@ -267,12 +305,14 @@ hip_add_exe_to_target(NAME ModuleTest
  TEST_SRC ${TEST_SRC}
  TEST_TARGET_NAME build_tests
  LINKER_LIBS ${RTCLIB}
-  COMMON_SHARED_SRC ${COMMON_SHARED_SRC}
-  COMPILE_OPTIONS -std=c++17)
+  COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+
+if(NOT WIN32)
+  add_dependencies(ModuleTest coopKernel.code)
+  add_dependencies(ModuleTest launch_kernel_module)
+endif()

-add_dependencies(ModuleTest coopKernel.code)
 add_dependencies(ModuleTest get_function_module)
-add_dependencies(ModuleTest launch_kernel_module)
 add_dependencies(ModuleTest get_global_test_module)
 add_dependencies(ModuleTest get_tex_ref_module)
 add_dependencies(ModuleTest not_a_module)
@@ -282,8 +322,11 @@ add_dependencies(ModuleTest empty_module)
 if(HIP_PLATFORM MATCHES "amd")
 add_dependencies(build_tests copyKernel.code copyKernel.s)
 add_dependencies(build_tests addKernel.code)
-add_dependencies(build_tests addKernel.spv)
-add_dependencies(build_tests addKernel-bundle.spv)
+# TODO - Rock build failure from clang
+if(NOT WIN32)
+  add_dependencies(build_tests addKernel.spv)
+  add_dependencies(build_tests addKernel-bundle.spv)
+endif()
 add_dependencies(build_tests copyKernelCompressed.code)
 add_dependencies(build_tests copyKernelGenericTarget.code)
 add_dependencies(build_tests copyKernelGenericTargetCompressed.code)
@@ -295,7 +338,10 @@ endif()
 endif()

 add_executable(hipGetFuncBySymbol_exe EXCLUDE_FROM_ALL hipGetFuncBySymbol_exe.cc)
+set_source_files_properties(hipGetFuncBySymbol_exe.cc PROPERTIES LANGUAGE HIP)
+set_target_properties(hipGetFuncBySymbol_exe PROPERTIES LINKER_LANGUAGE HIP)
 add_dependencies(build_tests hipGetFuncBySymbol_exe)
+target_link_libraries(hipGetFuncBySymbol_exe hip::host hip::device)
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS hipGetFuncBySymbol_exe)

 # Common Tests - Test independent of all platforms
@@ -318,15 +364,35 @@ endif()
 hip_add_exe_to_target(NAME module
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests)
-add_custom_target(managed_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/managed_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/managed_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(managed_kernel.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/managed_kernel.cpp
+  -o ${CMAKE_CURRENT_BINARY_DIR}/../module/managed_kernel.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(managed_kernel.code hip::host hip::device)

-add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cpp
+  -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(vcpy_kernel.code hip::host hip::device)

-add_custom_target(matmul.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/matmul.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/matmul.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(matmul.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  -x hip ${OFFLOAD_ARCH_LIST} ${CMAKE_CURRENT_SOURCE_DIR}/matmul.cpp -o
+  ${CMAKE_CURRENT_BINARY_DIR}/../module/matmul.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(matmul.code hip::host hip::device)

-add_custom_target(kernel_count.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_count.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_count.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(kernel_count.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/kernel_count.cpp -o
+  ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_count.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(kernel_count.code hip::host hip::device)

-add_custom_target(emptyModuleCount.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/emptyModuleCount.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/emptyModuleCount.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(emptyModuleCount.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/emptyModuleCount.cpp -o
+  ${CMAKE_CURRENT_BINARY_DIR}/../module/emptyModuleCount.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(emptyModuleCount.code hip::host hip::device)

 add_dependencies(ModuleTest managed_kernel.code)
 add_dependencies(ModuleTest vcpy_kernel.code)
@@ -334,7 +400,11 @@ add_dependencies(ModuleTest matmul.code)
 add_dependencies(ModuleTest kernel_count.code)
 add_dependencies(ModuleTest emptyModuleCount.code)

-add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH})
+add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_HIP_COMPILER} --cuda-device-only
+  ${OFFLOAD_ARCH_LIST} -x hip ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cpp -o
+  ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code
+  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include ${HIP_PATH_OPT})
+#target_link_libraries(kernel_composite_test.code hip::host hip::device)
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/kernel_composite_test.code
  ${CMAKE_CURRENT_BINARY_DIR}/matmul.code
@@ -344,7 +414,10 @@ set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/emptyModuleCount.code
 )
 add_executable(testhipModuleLoadUnloadFunc_exe EXCLUDE_FROM_ALL testhipModuleLoadUnloadFunc_exe.cc)
+set_source_files_properties(testhipModuleLoadUnloadFunc_exe.cc PROPERTIES LANGUAGE HIP)
+set_target_properties(testhipModuleLoadUnloadFunc_exe PROPERTIES LINKER_LANGUAGE HIP)
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS testhipModuleLoadUnloadFunc_exe)
+target_link_libraries(testhipModuleLoadUnloadFunc_exe hip::host hip::device)

 add_dependencies(module managed_kernel.code vcpy_kernel.code matmul.code kernel_composite_test.code
 testhipModuleLoadUnloadFunc_exe)
@@ -18,7 +18,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
 #include <hip/hip_cooperative_groups.h>

 using namespace cooperative_groups;
@@ -29,7 +29,10 @@ __global__ void cooperativeKernelEx(int* output, int totalThreads) {
  if (tid < totalThreads) {
    output[tid] = tid * 3;
  }
-  grid.sync();
+  // TODO: remove syncthreads with grid sync when compiler issue is fixed
+  // grid.sync();
+  __syncthreads();
+
  if (tid == 0) {
    output[0] = 2222;
  }
@@ -51,7 +54,9 @@ __global__ void argKernel(int* val) { *val = 100; }
 */
 __global__ void coopEmptykernel() {
  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
-  grid.sync();
+  // TODO: remove syncthreads with grid sync when compiler issue is fixed
+  // grid.sync();
+  __syncthreads();
 }

 /*
@@ -86,7 +91,9 @@ __global__ void coopFillArrayKernel(int* arr, int* output, int N) {
  else if (blockIdx.x == 9)
    arr[9] = 100;

-  grid.sync();
+  // TODO: remove syncthreads with grid sync when compiler issue is fixed
+  // grid.sync();
+  __syncthreads();

  for (int i = 0; i < N; i++) {
    output[blockIdx.x] = output[blockIdx.x] + arr[i];
@@ -17,7 +17,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>

 extern "C" __global__ void kernelMultipleArgsSaxpy(int a1, int a2, int* x1, int b1, int b2, int* x2,
                                                   int c1, int c2, int* x3, int d1, int d2, int* x4,
@@ -19,7 +19,7 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>

 extern "C" __global__ void copy_ker(int* Ad, int* Bd, size_t size) {
  int myId = threadIdx.x + blockDim.x * blockIdx.x;
@@ -19,7 +19,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
 #include "hip/hip_runtime_api.h"

 #include "hipModuleGetGlobal.hh"
@@ -113,7 +113,7 @@ TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Functional", "[multigpu]") {
    launchParamsList[i].args = args + i * NUM_KERNEL_ARGS;
  }

-  INFO("info: launch vector_square kernel with")
+  INFO("info: launch vector_square kernel with");
  INFO("hipExtLaunchMultiKernelMultiDevice API\n");
  HIP_CHECK(hipExtLaunchMultiKernelMultiDevice(launchParamsList, nGpu, 0));

@@ -17,7 +17,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>

 #define ARR_SIZE (32 * 32)
 #define SIZE (ARR_SIZE * sizeof(int))
@@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
 constexpr int GLOBAL_BUF_SIZE = 2048;

 __device__ float deviceGlobalFloat;
@@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>

 extern "C" __device__ void hello_world_2(float* a, float* b) {
    int tx = threadIdx.x;
@@ -38,6 +38,8 @@ __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {

 __global__ void CoopKernel() {
  cooperative_groups::grid_group grid = cooperative_groups::this_grid();
-  grid.sync();
+  // TODO: remove syncthreads with grid sync when compiler issue is fixed
+  // grid.sync();
+  __syncthreads();
 }
 }
@@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
 __managed__ int x = 10;

 extern "C" __global__ void GPU_func() { x++; }
@@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>
 __device__ int deviceGlobal = 1;

 extern "C" __global__ void matmulK(int clockrate, int* A, int* B, int* C, int N) {
@@ -16,7 +16,7 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include "hip/hip_runtime.h"
+#include <hip/hip_runtime.h>

 extern "C" __global__ void hello_world(float* a, float* b) {
  int tx = threadIdx.x;