[SYCL] Add Noexcept to clang SPIR-V builtin declaration (#17435)

wenju-he · web-flow · commit a20c57dcd526 · 2025-03-20T06:50:08.000-07:00
This fixes following error when SYCL header declares a SPIR-V builtin
with noexcept specifier and the builtin is also declared in
clang/lib/Sema/SPIRVBuiltins.td:
"exception specification in declaration does not match previous
declaration" For instance, __spirv_ControlBarrier declaration in
clang/test/CodeGenSYCL/Inputs/sycl.hpp has the build error.
diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp
@@ -786,6 +786,7 @@ static void GetProgModelBuiltinFctOverloads(
   FunctionProtoType::ExtProtoInfo PI(
       Context.getDefaultCallingConvention(false, false, true));
   PI.Variadic = IsVariadic;
+  PI.ExceptionSpec = FunctionProtoType::ExceptionSpecInfo{EST_BasicNoexcept};
 
   // Do not attempt to create any FunctionTypes if there are no return types,
   // which happens when a type belongs to a disabled extension.
diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-noexcept.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-noexcept.cpp
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fdeclare-spirv-builtins -fsyntax-only %s -verify
+// expected-no-diagnostics
+
+void acos(float val) {
+  static_assert(noexcept(__spirv_ocl_acos(val)));
+}
+
+void isnan(float a) {
+  static_assert(noexcept(__spirv_IsNan(a)));
+}
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -19,7 +19,7 @@ extern "C" int printf(const char* fmt, ...);
 
 #ifdef __SYCL_DEVICE_ONLY__
 __attribute__((convergent)) extern __attribute__((sycl_device)) void
-__spirv_ControlBarrier(int, int, int);
+__spirv_ControlBarrier(int, int, int) noexcept;
 #endif
 
 // Dummy runtime classes to model SYCL API.
diff --git a/libdevice/include/sanitizer_defs.hpp b/libdevice/include/sanitizer_defs.hpp
@@ -42,14 +42,15 @@ extern SYCL_EXTERNAL int
 __spirv_ocl_printf(const __SYCL_CONSTANT__ char *Format, ...);
 
 extern SYCL_EXTERNAL __SYCL_GLOBAL__ void *
-__spirv_GenericCastToPtrExplicit_ToGlobal(void *, int);
+__spirv_GenericCastToPtrExplicit_ToGlobal(void *, int) noexcept;
 extern SYCL_EXTERNAL __SYCL_LOCAL__ void *
-__spirv_GenericCastToPtrExplicit_ToLocal(void *, int);
+__spirv_GenericCastToPtrExplicit_ToLocal(void *, int) noexcept;
 extern SYCL_EXTERNAL __SYCL_PRIVATE__ void *
-__spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
+__spirv_GenericCastToPtrExplicit_ToPrivate(void *, int) noexcept;
 
 extern SYCL_EXTERNAL __attribute__((convergent)) void
-__spirv_ControlBarrier(int32_t Execution, int32_t Memory, int32_t Semantics);
+__spirv_ControlBarrier(int32_t Execution, int32_t Memory,
+                       int32_t Semantics) noexcept;
 
 extern "C" SYCL_EXTERNAL void __devicelib_exit();
 
diff --git a/libdevice/nativecpu_utils.cpp b/libdevice/nativecpu_utils.cpp
@@ -34,17 +34,19 @@ using __nativecpu_state = native_cpu::state;
 #define OCL_PRIVATE __attribute__((opencl_private))
 
 DEVICE_EXTERN_C void __mux_work_group_barrier(int32_t id, int32_t scope,
-                                              int32_t semantics);
+                                              int32_t semantics) noexcept;
 __SYCL_CONVERGENT__ DEVICE_EXTERNAL void
-__spirv_ControlBarrier(int32_t Execution, int32_t Memory, int32_t Semantics) {
+__spirv_ControlBarrier(int32_t Execution, int32_t Memory,
+                       int32_t Semantics) noexcept {
   if (__spv::Scope::Flag::Workgroup == Execution)
     // todo: check id and args; use mux constants
     __mux_work_group_barrier(0, Execution, Semantics);
 }
 
-DEVICE_EXTERN_C void __mux_mem_barrier(int32_t scope, int32_t semantics);
+DEVICE_EXTERN_C void __mux_mem_barrier(int32_t scope,
+                                       int32_t semantics) noexcept;
 __SYCL_CONVERGENT__ DEVICE_EXTERNAL void
-__spirv_MemoryBarrier(int32_t Memory, int32_t Semantics) {
+__spirv_MemoryBarrier(int32_t Memory, int32_t Semantics) noexcept {
   __mux_mem_barrier(Memory, Semantics);
 }
 
@@ -54,7 +56,7 @@ __spirv_MemoryBarrier(int32_t Memory, int32_t Semantics) {
 
 #define DefGenericCastToPtrExplImpl(sfx, asp, cv)\
 DEVICE_EXTERNAL cv asp void *\
-__spirv_GenericCastToPtrExplicit_##sfx(cv void *p ,int) {\
+__spirv_GenericCastToPtrExplicit_##sfx(cv void *p ,int) noexcept {\
   return (cv asp void *)p;\
 }
 
@@ -100,9 +102,9 @@ DefSubgroupBlockINTEL(uint32_t) DefSubgroupBlockINTEL(uint64_t)
 DefSubgroupBlockINTEL(uint8_t) DefSubgroupBlockINTEL(uint16_t)
 
 #define DefineGOp1(spir_sfx, name)\
-DEVICE_EXTERN_C bool __mux_sub_group_##name##_i1(bool);\
-DEVICE_EXTERN_C bool __mux_work_group_##name##_i1(uint32_t id, bool val);\
-DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(unsigned g, bool val) {\
+DEVICE_EXTERN_C bool __mux_sub_group_##name##_i1(bool) noexcept;\
+DEVICE_EXTERN_C bool __mux_work_group_##name##_i1(uint32_t id, bool val) noexcept;\
+DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(unsigned g, bool val) noexcept {\
   if (__spv::Scope::Flag::Subgroup == g)\
     return __mux_sub_group_##name##_i1(val);\
   else if (__spv::Scope::Flag::Workgroup == g)\
@@ -115,16 +117,16 @@ DefineGOp1(All, all)
 
 
 #define DefineGOp(Type, MuxType, spir_sfx, mux_sfx)                            \
-  DEVICE_EXTERN_C MuxType __mux_sub_group_scan_inclusive_##mux_sfx(MuxType);   \
-  DEVICE_EXTERN_C MuxType __mux_sub_group_scan_exclusive_##mux_sfx(MuxType);   \
-  DEVICE_EXTERN_C MuxType __mux_sub_group_reduce_##mux_sfx(MuxType);           \
+  DEVICE_EXTERN_C MuxType __mux_sub_group_scan_inclusive_##mux_sfx(MuxType) noexcept;   \
+  DEVICE_EXTERN_C MuxType __mux_sub_group_scan_exclusive_##mux_sfx(MuxType) noexcept;   \
+  DEVICE_EXTERN_C MuxType __mux_sub_group_reduce_##mux_sfx(MuxType) noexcept;           \
   DEVICE_EXTERN_C MuxType __mux_work_group_scan_exclusive_##mux_sfx(uint32_t,  \
-                                                                    MuxType);  \
+                                                                    MuxType) noexcept;  \
   DEVICE_EXTERN_C MuxType __mux_work_group_scan_inclusive_##mux_sfx(uint32_t,  \
-                                                                    MuxType);  \
-  DEVICE_EXTERN_C MuxType __mux_work_group_reduce_##mux_sfx(uint32_t, MuxType);\
+                                                                    MuxType) noexcept;  \
+  DEVICE_EXTERN_C MuxType __mux_work_group_reduce_##mux_sfx(uint32_t, MuxType) noexcept;\
   DEVICE_EXTERNAL Type __spirv_Group##spir_sfx(uint32_t g, uint32_t id,        \
-                                               Type v) {                       \
+                                               Type v) noexcept {              \
     if (__spv::Scope::Flag::Subgroup == g) {                                   \
       if (static_cast<unsigned>(__spv::GroupOperation::InclusiveScan) == id)   \
         return __mux_sub_group_scan_inclusive_##mux_sfx(v);                    \
@@ -196,29 +198,29 @@ DefineLogicalGroupOp(bool, bool, i1)
 
 #define DefineBroadcastMuxType(Type, Sfx, MuxType, IDType)                    \
   DEVICE_EXTERN_C MuxType __mux_work_group_broadcast_##Sfx(                   \
-      int32_t id, MuxType val, uint64_t lidx, uint64_t lidy, uint64_t lidz);  \
+      int32_t id, MuxType val, uint64_t lidx, uint64_t lidy, uint64_t lidz) noexcept; \
   DEVICE_EXTERN_C MuxType __mux_sub_group_broadcast_##Sfx(MuxType val,        \
-                                                          int32_t sg_lid);
+                                                          int32_t sg_lid) noexcept;
 
 #define DefineBroadCastImpl(Type, Sfx, MuxType, IDType)                       \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v,             \
-                                              IDType l) {                     \
+                                              IDType l) noexcept {            \
     if (__spv::Scope::Flag::Subgroup == g)                                    \
       return __mux_sub_group_broadcast_##Sfx(v, l);                           \
     else                                                                      \
       return __mux_work_group_broadcast_##Sfx(0, v, l, 0, 0);                 \
   }                                                                           \
                                                                               \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v,             \
-                                         sycl::vec<IDType, 2>::vector_t l) {  \
+                                  sycl::vec<IDType, 2>::vector_t l) noexcept{ \
     if (__spv::Scope::Flag::Subgroup == g)                                    \
       return __mux_sub_group_broadcast_##Sfx(v, l[0]);                        \
     else                                                                      \
       return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[1], 0);           \
   }                                                                           \
                                                                               \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v,             \
-                                          sycl::vec<IDType, 3>::vector_t l) { \
+                                  sycl::vec<IDType, 3>::vector_t l) noexcept{ \
     if (__spv::Scope::Flag::Subgroup == g)                                    \
       return __mux_sub_group_broadcast_##Sfx(v, l[0]);                        \
     else                                                                      \
@@ -241,7 +243,7 @@ DefineBroadCast(int64_t, i64, int64_t)
 
 #define DefShuffleINTEL(Type, Sfx, MuxType)                                    \
   DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_##Sfx(MuxType val,           \
-                                                        int32_t lid);          \
+                                                        int32_t lid) noexcept; \
   template <>                                                                  \
   DEVICE_EXTERNAL Type __spirv_SubgroupShuffleINTEL<Type>(                     \
       Type val, unsigned id) noexcept {                                        \
@@ -250,7 +252,7 @@ DefineBroadCast(int64_t, i64, int64_t)
 
 #define DefShuffleUpINTEL(Type, Sfx, MuxType)                                  \
   DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_up_##Sfx(                    \
-      MuxType prev, MuxType curr, int32_t delta);                              \
+      MuxType prev, MuxType curr, int32_t delta) noexcept;                     \
   template <>                                                                  \
   DEVICE_EXTERNAL Type __spirv_SubgroupShuffleUpINTEL<Type>(                   \
       Type prev, Type curr, unsigned delta) noexcept {                         \
@@ -260,7 +262,7 @@ DefineBroadCast(int64_t, i64, int64_t)
 
 #define DefShuffleDownINTEL(Type, Sfx, MuxType)                                \
   DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_down_##Sfx(                  \
-      MuxType curr, MuxType next, int32_t delta);                              \
+      MuxType curr, MuxType next, int32_t delta) noexcept;                     \
   template <>                                                                  \
   DEVICE_EXTERNAL Type __spirv_SubgroupShuffleDownINTEL<Type>(                 \
       Type curr, Type next, unsigned delta) noexcept {                         \
@@ -298,7 +300,7 @@ DefShuffleINTEL_All(_Float16, f16, _Float16)
 // Vector versions of shuffle are generated by the FixABIBuiltinsSYCLNativeCPU pass
 
 #define Define2ArgForward(Type, Name, Callee)\
-DEVICE_EXTERNAL Type Name(Type a, Type b) { return Callee(a,b);}
+DEVICE_EXTERNAL Type Name(Type a, Type b) noexcept { return Callee(a,b);}
 
 Define2ArgForward(uint64_t, __spirv_ocl_u_min, std::min)
 
