merge amd-stg-open into amd-mainline-open

Change-Id: I26c76821066a0cd72ab13a9860befcaa997b1cc0
ROCm · Feb 16, 2022 · 11a6689 · 11a6689
2 parents 62809a5 + 4dfb181
commit 11a6689
Show file tree

Hide file tree

Showing 9 changed files with 56 additions and 10 deletions.
diff --git a/asanrtl/src/globals.cl b/asanrtl/src/globals.cl
@@ -45,7 +45,7 @@ poison_redzones(__global const struct device_global *g) {
 }
 
 // This function is called by one-workitem constructor kernel.
-NO_SANITIZE_ADDR
+OPT_NONE NO_SANITIZE_ADDR
 void
 __asan_register_globals(uptr globals, uptr n) {
     __global struct device_global *dglobals = (__global struct device_global*) globals;
@@ -65,7 +65,7 @@ unpoison_global(__global const struct device_global *g) {
 }
 
 // This function is called by one-workitem destructor kernel.
-NO_SANITIZE_ADDR
+OPT_NONE NO_SANITIZE_ADDR
 void
 __asan_unregister_globals(uptr globals, uptr n) {
     __global struct device_global* dglobals = (__global struct device_global*) globals;

diff --git a/asanrtl/src/preserve.cl b/asanrtl/src/preserve.cl
@@ -5,7 +5,7 @@
  * License. See LICENSE.TXT for details.
  *===------------------------------------------------------------------------*/
 
-typedef ulong uptr;
+#include "asan_util.h"
 extern void __asan_report_load1 (uptr addr);
 extern void __asan_report_load1_noabort (uptr addr);
 extern void __asan_report_load2 (uptr addr);
@@ -72,7 +72,10 @@ extern void __asan_unregister_elf_globals(uptr flag, uptr start, uptr stop);
 extern void __asan_init(void);
 extern void __asan_version_mismatch_check_v8(void);
 
-void
+//  Functions called within the below function must not get inlined and their
+//  names want to be preserved at higher opt level to enable linking between
+//  module compiled with asan instrumentation and asan device rtl.
+OPT_NONE void
 __amdgpu_device_library_preserve_asan_functions(void)
 {
     __asan_report_load1(0);

diff --git a/asanrtl/src/shadow_mapping.cl b/asanrtl/src/shadow_mapping.cl
@@ -30,7 +30,7 @@ range_check(uptr beg, uptr end) {
 }
 
 //check all application bytes in [beg,beg+size) range are accessible
-NO_SANITIZE_ADDR
+OPT_NONE NO_SANITIZE_ADDR
 uptr
 __asan_region_is_poisoned(uptr beg, uptr size)
 {

diff --git a/asanrtl/src/stubs.cl b/asanrtl/src/stubs.cl
@@ -5,7 +5,7 @@
  * License. See LICENSE.TXT for details.
  *===------------------------------------------------------------------------*/
 
-typedef ulong uptr;
+#include "asan_util.h"
 
 void __asan_handle_no_return(void) {}
 
@@ -25,7 +25,7 @@ void __asan_register_elf_globals(uptr flag, uptr start, uptr stop) {}
 
 void __asan_unregister_elf_globals(uptr flag, uptr start, uptr stop) {}
 
-void __asan_init(void) {}
+OPT_NONE NO_SANITIZE_ADDR void __asan_init(void) {}
 
-void __asan_version_mismatch_check_v8(void) {}
+OPT_NONE NO_SANITIZE_ADDR void __asan_version_mismatch_check_v8(void) {}
 
diff --git a/cmake/OCL.cmake b/cmake/OCL.cmake
@@ -29,6 +29,18 @@ if (WIN32)
   set(CLANG_OCL_FLAGS ${CLANG_OCL_FLAGS} -fshort-wchar)
 endif()
 
+# Disable code object version module flag if available.
+file(WRITE ${CMAKE_BINARY_DIR}/tmp.cl "")
+execute_process (
+  COMMAND ${LLVM_TOOLS_BINARY_DIR}/clang${EXE_SUFFIX} ${CLANG_OCL_FLAGS} -Xclang -mcode-object-version=none ${CMAKE_BINARY_DIR}/tmp.cl
+  RESULT_VARIABLE TEST_CODE_OBJECT_VERSION_NONE_RESULT
+  ERROR_QUIET
+)
+file(REMOVE ${CMAKE_BINARY_DIR}/tmp.cl)
+if (NOT TEST_CODE_OBJECT_VERSION_NONE_RESULT)
+  set(CLANG_OCL_FLAGS ${CLANG_OCL_FLAGS} -Xclang -mcode-object-version=none)
+endif()
+
 set (BC_EXT .bc)
 set (LIB_SUFFIX ".lib${BC_EXT}")
 set (STRIP_SUFFIX ".strip${BC_EXT}")

diff --git a/ockl/src/cg.cl b/ockl/src/cg.cl
@@ -82,6 +82,7 @@ __attribute__((convergent)) void
 __ockl_grid_sync(void)
 {
     __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+    __builtin_amdgcn_s_barrier();
     if (choose_one_workgroup_workitem()) {
         uint nwm1 = (uint)__ockl_get_num_groups(0) * (uint)__ockl_get_num_groups(1) * (uint)__ockl_get_num_groups(2) - 1;
         __ockl_gws_barrier(nwm1, 0);
@@ -125,14 +126,15 @@ __ockl_multi_grid_is_valid(void)
 __attribute__((convergent)) void
 __ockl_multi_grid_sync(void)
 {
-    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+    __builtin_amdgcn_s_barrier();
     uint nwm1 = (uint)__ockl_get_num_groups(0) * (uint)__ockl_get_num_groups(1) * (uint)__ockl_get_num_groups(2) - 1;
     bool cwwi = choose_one_workgroup_workitem();
 
     if (cwwi)
         __ockl_gws_barrier(nwm1, 0);
 
-    __builtin_amdgcn_s_barrier();
+    // Need another barrier here if the following choose doesn't see cwwi set
 
     if (choose_one_grid_workitem()) {
         __constant struct mg_info *m = (__constant struct mg_info *)get_mg_info_arg();

diff --git a/oclc/inc/oclc.h b/oclc/inc/oclc.h
@@ -25,9 +25,15 @@
 //    __constant bool __oclc_correctly_rounded_sqrt32(void)
 //        - the application is expecting sqrt(float) to produce a correctly rounded result
 //
+//    __constant bool __oclc_wavefrontsize64
+//        - the application is being compiled for a wavefront size of 64
+//
 //    __constant int __oclc_ISA_version
 //        - the ISA version of the target device
 //
+//    __constant int __oclc_ABI_version
+//        - the ABI version the application is being compiled for
+//
 // it is expected that the implementation provides these as if declared from the following
 // C code:
 //
@@ -41,5 +47,6 @@ extern const __constant bool __oclc_daz_opt;
 extern const __constant bool __oclc_correctly_rounded_sqrt32;
 extern const __constant bool __oclc_wavefrontsize64;
 extern const __constant int __oclc_ISA_version;
+extern const __constant int __oclc_ABI_version;
 
 #endif // OCLC_H
diff --git a/oclc/src/abi_version_400.cl b/oclc/src/abi_version_400.cl
@@ -0,0 +1,11 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ABI_version = 400;
+
diff --git a/oclc/src/abi_version_500.cl b/oclc/src/abi_version_500.cl
@@ -0,0 +1,11 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ABI_version = 500;
+