__HIP_DEVICE_COMPILE__ do_not_optimize (broken)

cwpearson · Nov 4, 2022 · 3392fd0 · 3392fd0
1 parent 210a2e6
commit 3392fd0
Show file tree

Hide file tree

Showing 2 changed files with 16 additions and 0 deletions.
diff --git a/include/scope/do_not_optimize.hpp b/include/scope/do_not_optimize.hpp
@@ -0,0 +1,15 @@
+#pragma once
+
+namespace scope {
+template <class Tp>
+__device__
+void __attribute__((always_inline)) do_not_optimize(Tp const& value) {
+  asm volatile("" : : "r,m"(value) : "memory");
+}
+
+template <class Tp>
+__device__
+void __attribute__((always_inline)) do_not_optimize(Tp& value) {
+  asm volatile("" : "+r,m"(value) : : "memory");
+}
+}
diff --git a/include/scope/scope.hpp b/include/scope/scope.hpp
@@ -16,6 +16,7 @@
 #if defined(SCOPE_USE_HIP)
 #include "scope/hip.hpp"
 #if defined(__HIP_DEVICE_COMPILE__)
+#include "scope/do_not_optimize.hpp"
 #endif // __HIP_DEVICE_COMPILE__
 #endif // SCOPE_USE_HIP