diff --git a/pips/device/device_common.hpp b/pips/device/device_common.hpp
index ef41ec5..641b3d5 100644
--- a/pips/device/device_common.hpp
+++ b/pips/device/device_common.hpp
@@ -1,17 +1,17 @@
 #ifndef PIPS_DEVICE_DEVICE_COMMON_HPP_
 #define PIPS_DEVICE_DEVICE_COMMON_HPP_
 
-// Qualifier used on functions that must be callable from both host code and
-// CUDA/HIP device code. On plain C++ compilers it expands to nothing.
+// Qualifiers for functions that must be callable from both host code and
+// CUDA/HIP device code.
 #if defined(__CUDACC__) || defined(__HIPCC__)
-#define PIPS_DEVICE_HOST __host__ __device__ inline
-#define PIPS_DEVICE_HOST_INLINE PIPS_DEVICE_HOST __forceinline__
+#define PIPS_DEVICE_HOST_INLINE __host__ __device__ inline
+#define PIPS_DEVICE_HOST_FORCEINLINE __host__ __device__ __forceinline__
 #else
-#define PIPS_DEVICE_HOST inline
+#define PIPS_DEVICE_HOST_INLINE inline
 #if defined(__GNUC__) || defined(__clang__)
-#define PIPS_DEVICE_HOST_INLINE __attribute__((always_inline)) inline
+#define PIPS_DEVICE_HOST_FORCEINLINE __attribute__((always_inline)) inline
 #else
-#define PIPS_DEVICE_HOST_INLINE PIPS_DEVICE_HOST
+#define PIPS_DEVICE_HOST_FORCEINLINE inline
 #endif
 #endif
 
diff --git a/pips/device/device_value.hpp b/pips/device/device_value.hpp
index 1d0c16b..a11c1f4 100644
--- a/pips/device/device_value.hpp
+++ b/pips/device/device_value.hpp
@@ -45,28 +45,28 @@ struct DeviceValue {
   } as;
 };
 
-PIPS_DEVICE_HOST_INLINE DeviceValue dv_nil() {
+PIPS_DEVICE_HOST_FORCEINLINE DeviceValue dv_nil() {
   DeviceValue v{};
   v.type = DeviceValueType::NIL;
   v.as.n = 0;
   return v;
 }
 
-PIPS_DEVICE_HOST_INLINE DeviceValue dv_bool(bool x) {
+PIPS_DEVICE_HOST_FORCEINLINE DeviceValue dv_bool(bool x) {
   DeviceValue v{};
   v.type = DeviceValueType::BOOL;
   v.as.b = x;
   return v;
 }
 
-PIPS_DEVICE_HOST_INLINE DeviceValue dv_number(DeviceReal x) {
+PIPS_DEVICE_HOST_FORCEINLINE DeviceValue dv_number(DeviceReal x) {
   DeviceValue v{};
   v.type = DeviceValueType::NUMBER;
   v.as.n = x;
   return v;
 }
 
-PIPS_DEVICE_HOST_INLINE bool dv_vector(const DeviceReal *elements,
+PIPS_DEVICE_HOST_FORCEINLINE bool dv_vector(const DeviceReal *elements,
                                        std::uint32_t length,
                                        DeviceValue &out) {
   if (length > PIPS_DEVICE_VECTOR_MAX || (length > 0 && !elements))
@@ -80,37 +80,37 @@ PIPS_DEVICE_HOST_INLINE bool dv_vector(const DeviceReal *elements,
   return true;
 }
 
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_is_nil(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_is_nil(const DeviceValue &v) {
   return v.type == DeviceValueType::NIL;
 }
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_is_bool(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_is_bool(const DeviceValue &v) {
   return v.type == DeviceValueType::BOOL;
 }
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_is_number(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_is_number(const DeviceValue &v) {
   return v.type == DeviceValueType::NUMBER;
 }
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_is_vector(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_is_vector(const DeviceValue &v) {
   return v.type == DeviceValueType::VECTOR;
 }
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_as_bool(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_as_bool(const DeviceValue &v) {
   return v.as.b;
 }
-PIPS_DEVICE_HOST_INLINE constexpr DeviceReal dv_as_number(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr DeviceReal dv_as_number(const DeviceValue &v) {
   return v.as.n;
 }
-PIPS_DEVICE_HOST_INLINE constexpr std::uint8_t
+PIPS_DEVICE_HOST_FORCEINLINE constexpr std::uint8_t
 dv_vector_length(const DeviceValue &v) {
   return v.as.vector.length;
 }
-PIPS_DEVICE_HOST_INLINE constexpr DeviceReal
+PIPS_DEVICE_HOST_FORCEINLINE constexpr DeviceReal
 dv_vector_element(const DeviceValue &v, std::uint8_t index) {
   return v.as.vector.elements[index];
 }
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_vector_is_valid(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_vector_is_valid(const DeviceValue &v) {
   return !dv_is_vector(v) || v.as.vector.length <= PIPS_DEVICE_VECTOR_MAX;
 }
 
-PIPS_DEVICE_HOST_INLINE constexpr bool dv_is_falsey(const DeviceValue &v) {
+PIPS_DEVICE_HOST_FORCEINLINE constexpr bool dv_is_falsey(const DeviceValue &v) {
   return dv_is_nil(v) || (dv_is_bool(v) && !v.as.b);
 }
 
diff --git a/pips/device/device_vm.hpp b/pips/device/device_vm.hpp
index 3870317..fb6b584 100644
--- a/pips/device/device_vm.hpp
+++ b/pips/device/device_vm.hpp
@@ -34,18 +34,18 @@ struct DeviceVM {
   std::int32_t sp = 0; 
   std::int32_t fp = 0; 
 
-  PIPS_DEVICE_HOST_INLINE bool push(const DeviceValue &v) {
+  PIPS_DEVICE_HOST_FORCEINLINE bool push(const DeviceValue &v) {
     if (sp >= PIPS_DEVICE_STACK_MAX) return false;
     stack[sp++] = v;
     return true;
   }
-  PIPS_DEVICE_HOST_INLINE DeviceValue pop() { return stack[--sp]; }
-  PIPS_DEVICE_HOST_INLINE const DeviceValue &peek(int back) const {
+  PIPS_DEVICE_HOST_FORCEINLINE DeviceValue pop() { return stack[--sp]; }
+  PIPS_DEVICE_HOST_FORCEINLINE const DeviceValue &peek(int back) const {
     return stack[sp - 1 - back];
   }
 
   // Execute a packed function
-  PIPS_DEVICE_HOST DeviceStatus run(const DeviceModule &module,
+  PIPS_DEVICE_HOST_INLINE DeviceStatus run(const DeviceModule &module,
                                     std::uint32_t entry_id,
                                     const DeviceValue *args,
                                     std::uint32_t argc,
@@ -72,7 +72,7 @@ struct DeviceVM {
   }
 
 private:
-  PIPS_DEVICE_HOST_INLINE static DeviceStatus
+  PIPS_DEVICE_HOST_FORCEINLINE static DeviceStatus
   scalar_arith(DeviceOpCode op, DeviceReal a, DeviceReal b, DeviceReal &out) {
     switch (op) {
     case DeviceOpCode::ADD: out = a + b; return DeviceStatus::OK;
@@ -93,7 +93,7 @@ struct DeviceVM {
     }
   }
 
-  PIPS_DEVICE_HOST_INLINE static DeviceStatus
+  PIPS_DEVICE_HOST_FORCEINLINE static DeviceStatus
   vector_arith(DeviceOpCode op, const DeviceValue &a, const DeviceValue &b,
                DeviceValue &out) {
     const bool av = dv_is_vector(a);
@@ -128,7 +128,7 @@ struct DeviceVM {
     return DeviceStatus::OK;
   }
 
-  PIPS_DEVICE_HOST DeviceStatus dispatch(const DeviceModule &module,
+  PIPS_DEVICE_HOST_INLINE DeviceStatus dispatch(const DeviceModule &module,
                                          DeviceValue *out_result) {
     using OC = DeviceOpCode;