Re impl assert

libdevice/cmake/modules/SYCLLibdevice.cmake

@@ -2,11 +2,32 @@ set(obj_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}")
 if (WIN32)
   set(lib-suffix obj)
   set(spv_binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}")
+  set(host_offload_target "host-x86_64-pc-windows-msvc")
 else()
   set(lib-suffix o)
   set(spv_binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}")
+  set(host_offload_target "host-x86_64-unknown-linux-gnu")
 endif()
 set(clang $<TARGET_FILE:clang>)
+set(bundler $<TARGET_FILE:clang-offload-bundler>)
+set(llvm-link $<TARGET_FILE:llvm-link>)
+set(llvm-spirv $<TARGET_FILE:llvm-spirv>)
+
+string(CONCAT bundler_targets_opt
+  "-targets="
+  "sycl-spir64_x86_64-unknown-unknown,"
+  "sycl-spir64_gen-unknown-unknown,"
+  "sycl-spir64_fpga-unknown-unknown,"
+  "sycl-spir64-unknown-unknown,"
+  ${host_offload_target})
+string(CONCAT bundler_inputs_opt
+  "-inputs="
+  ${obj_binary_dir} "/fallback-cassert_spir64_x86_64." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64_gen." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64_fpga." ${lib-suffix} ","
+  ${obj_binary_dir} "/fallback-cassert_spir64." ${lib-suffix} ","
+  ${obj_binary_dir} "/assert_no_read_host." ${lib-suffix})
+
 
 string(CONCAT sycl_targets_opt
   "-fsycl-targets="
@@ -81,6 +102,61 @@ add_custom_command(OUTPUT ${devicelib-obj-cmath-fp64}
                    DEPENDS device_math.h device.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+#[[
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                   MAIN_DEPENDENCY fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   COMMAND ${llvm-link} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                           ${spv_binary_dir}/assert_read_spir64_spv.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert_no_read.bc
+                           ${spv_binary_dir}/assert_read_spir64_spv.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                   COMMAND ${llvm-spirv} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           llvm-spirv
+                   VERBATIM)
+]]
+
+
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   MAIN_DEPENDENCY fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                   COMMAND ${llvm-spirv} -o
+                           ${spv_binary_dir}/libsycl-fallback-cassert.spv
+                           ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                   DEPENDS ${spv_binary_dir}/libsycl-fallback-cassert.bc
+                           llvm-spirv
+                   VERBATIM)
+#[[
 add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
                    COMMAND ${clang} -fsycl-device-only -fno-sycl-use-bitcode
                            ${compile_opts}
@@ -89,6 +165,8 @@ add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cassert.spv
                    MAIN_DEPENDENCY fallback-cassert.cpp
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
+]]
+################################################################################
 
 add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cstring.spv
                    COMMAND ${clang} -fsycl-device-only -fno-sycl-use-bitcode
@@ -99,6 +177,127 @@ add_custom_command(OUTPUT ${spv_binary_dir}/libsycl-fallback-cstring.spv
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+#[[
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_x86_64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_gen.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_gen-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_gen.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64_fpga-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_read_spir64.bc
+                   COMMAND ${clang} -c -x cl -emit-llvm -fsycl-device-only
+                           --target=spir64-unknown-unknown -cl-std=CL2.0
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                           -o ${obj_binary_dir}/assert_read_spir64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cl
+                   DEPENDS sycl-compiler
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_x86_64-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_gen-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   COMMAND ${clang} -fsycl-device-only -emit-llvm
+                           -fsycl-targets=spir64_fpga-unknown-unknown ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/assert_no_read_host.o
+                   COMMAND ${clang} -c
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${obj_binary_dir}/assert_no_read_host.o
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                   DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_x86_64.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_x86_64.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_gen.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_gen.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_gen.bc
+                           llvm-link
+                   VERBATIM)
+add_custom_command(OUTPUT ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                   COMMAND ${llvm-link} -o ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                           ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                   DEPENDS ${obj_binary_dir}/assert_read_spir64_fpga.bc
+                           ${obj_binary_dir}/assert_no_read_spir64_fpga.bc
+                           llvm-link
+                   VERBATIM)
+
+add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
+                   COMMAND ${bundler} -type=o ${bundler_targets_opt}
+                           -outputs=${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
+                           ${bundler_inputs_opt}
+                   DEPENDS ${obj_binary_dir}/fallback-cassert_spir64_x86_64.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64_gen.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64_fpga.${lib-suffix}
+                           ${obj_binary_dir}/fallback-cassert_spir64.${lib-suffix}
+                           ${obj_binary_dir}/assert_no_read_host.${lib-suffix}
+                   VERBATIM)
+]]
+
+
 add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffix}
                    COMMAND ${clang} -fsycl -c
                            ${compile_opts} ${sycl_targets_opt}
@@ -108,6 +307,8 @@ add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cassert.${lib-suffi
                    DEPENDS wrapper.h device.h spirv_vars.h sycl-compiler
                    VERBATIM)
 
+################################################################################
+
 add_custom_command(OUTPUT ${obj_binary_dir}/libsycl-fallback-cstring.${lib-suffix}
                    COMMAND ${clang} -fsycl -c
                            ${compile_opts} ${sycl_targets_opt}

libdevice/fallback-cassert.cl

@@ -0,0 +1,51 @@
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics:enable
+#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics:enable
+
+// NOTE Align these definitions with fallback-cassert.cpp
+#define ASSERT_NONE 0
+#define ASSERT_START 1
+#define ASSERT_FINISH 2
+
+// NOTE Layout of this structure should be aligned with the one in
+// sycl/include/CL/sycl/detail/assert_happened.hpp
+struct AssertHappened {
+  int Flag;
+  char Expr[256 + 1];
+  char File[256 + 1];
+  char Func[128 + 1];
+
+  int Line;
+
+  unsigned long GID0;
+  unsigned long GID1;
+  unsigned long GID2;
+
+  unsigned long LID0;
+  unsigned long LID1;
+  unsigned long LID2;
+};
+
+typedef struct AssertHappened AssertHappenedT;
+
+extern __global AssertHappenedT SPIR_AssertHappenedMem;
+
+__kernel void __devicelib_assert_read(__global void *_Dst) {
+  if (!_Dst)
+    return;
+
+  AssertHappenedT *Dst = (AssertHappenedT *)_Dst;
+
+  __global int *FlagPtr = &SPIR_AssertHappenedMem.Flag;
+  int Flag = atomic_add(FlagPtr, 0);
+
+  if (ASSERT_NONE == Flag) {
+    Dst->Flag = ASSERT_NONE;
+    return;
+  }
+
+  if (Flag != ASSERT_FINISH)
+    while (ASSERT_START == atomic_add(FlagPtr, 0))
+      ;
+
+  *Dst = SPIR_AssertHappenedMem;
+}

libdevice/fallback-cassert.cpp

@@ -12,14 +12,15 @@
 
 #ifdef __SPIR__
 
+// NOTE Align these definitions with fallback-cassert.cl
 #define ASSERT_NONE 0
 #define ASSERT_START 1
 #define ASSERT_FINISH 2
 
 // definition
 SPIR_GLOBAL AssertHappened SPIR_AssertHappenedMem;
 
-DEVICE_EXTERN_C void __devicelib_assert_read(void *_Dst) {
+void __devicelib_assert_read(__SYCL_GLOBAL__ void *_Dst) {
   AssertHappened *Dst = (AssertHappened *)_Dst;
   int Flag = atomicLoad(&SPIR_AssertHappenedMem.Flag);
 
@@ -35,11 +36,58 @@ DEVICE_EXTERN_C void __devicelib_assert_read(void *_Dst) {
   *Dst = SPIR_AssertHappenedMem;
 }
 
+#ifdef SYCL_LANGUAGE_VERSION
+#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]]
+#else
+#define __SYCL_KERNEL_ATTR__
+#endif
+
+// Create a kernel entry point for __devicelib_assert_read
+template <typename KernelName, typename KernelType>
+__SYCL_KERNEL_ATTR__ __attribute__((noinline)) void
+kernel_caller(KernelType KernelFunc) {
+  KernelFunc();
+}
+
+namespace detail {
+class DevicelibAssertReadKernel {
+  __SYCL_GLOBAL__ void *MDst;
+public:
+  DevicelibAssertReadKernel(__SYCL_GLOBAL__ void *Dst) : MDst{Dst} {}
+  __attribute__((noinline))
+  void operator()() const {
+    __devicelib_assert_read(MDst);
+  }
+};
+} // namespace detail
+
+
+using DevicelibAssertReadT = detail::DevicelibAssertReadKernel;//void (*)(__SYCL_GLOBAL__ void *);
+
+__attribute__((noinline))
+struct DevicelibAssertReadKernelName;
+
+// a stub function to enforce assert info reader kernel in devicelib image
+__attribute__((noinline))
+void __devicelib_stub() {
+  kernel_caller<DevicelibAssertReadKernelName, DevicelibAssertReadT>(DevicelibAssertReadT{NULL});
+}
+
 DEVICE_EXTERN_C void __devicelib_assert_fail(const char *expr, const char *file,
                                              int32_t line, const char *func,
                                              uint64_t gid0, uint64_t gid1,
                                              uint64_t gid2, uint64_t lid0,
                                              uint64_t lid1, uint64_t lid2) {
+  // FIXME make offline linking against __devicelib_assert_fail enforce linking
+  // against __devicelib_assert_read also
+  {
+    __devicelib_stub();
+//    kernel_caller<DevicelibAssertReadKernelName, DevicelibAssertReadT>(DevicelibAssertReadT{NULL});
+//    __devicelib_stub();
+//    __devicelib_assert_read(NULL);
+  }
+
+
   int Expected = ASSERT_NONE;
   int Desired = ASSERT_START;
 

libdevice/include/assert-happened.hpp

@@ -43,6 +43,6 @@ struct AssertHappened {
 #define __SYCL_GLOBAL__ __attribute__((opencl_global))
 
 // declaration
-extern SPIR_GLOBAL_VAR __SYCL_GLOBAL__ AssertHappened SPIR_AssertHappenedMem;
+extern "C" SPIR_GLOBAL_VAR __SYCL_GLOBAL__ AssertHappened SPIR_AssertHappenedMem;
 
 #endif

sycl/include/CL/sycl/detail/assert_happened.hpp

@@ -12,12 +12,6 @@
 
 #include <cstdint>
 
-#ifdef __SYCL_DEVICE_ONLY__
-// Reads Flag of AssertHappened on device
-SYCL_EXTERNAL __attribute__((weak)) extern "C" void
-__devicelib_assert_read(void *);
-#endif
-
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {