Fix __builtin_amdgcn_s_sleep() not present on RustiCL (see preda#276)

olympichek · Dec 23, 2024 · e52572e · e52572e
1 parent b947294
commit e52572e
Showing 1 changed file with 3 additions and 1 deletion.
diff --git a/src/cl/carryfused.cl b/src/cl/carryfused.cl
@@ -6,8 +6,10 @@
 #include "middle.cl"
 
 void spin() {
-#if AMDGPU
+#if defined(__has_builtin) && __has_builtin(__builtin_amdgcn_s_sleep)
   __builtin_amdgcn_s_sleep(0);
+#elif HAS_ASM
+  __asm("s_sleep 0");
 #else
   // nothing: just spin
   // on Nvidia: see if there's some brief sleep function