diff --git a/configure.ac b/configure.ac
index 0f1c4fc..a79aed6 100644
--- a/configure.ac
+++ b/configure.ac
@@ -32,6 +32,7 @@ LT_INIT
 
 AC_SUBST(VERSION_INFO, [version_info] )
 
+m4_include([m4/ax_check_compile_flag.m4])
 AX_CHECK_COMPILE_FLAG([-std=c++0x], [CXXFLAGS="$CXXFLAGS -std=c++0x"])
 AX_PTHREAD
 
diff --git a/include/gdsync/device.cuh b/include/gdsync/device.cuh
index 25f5dfd..420d039 100644
--- a/include/gdsync/device.cuh
+++ b/include/gdsync/device.cuh
@@ -89,18 +89,6 @@ namespace gdsync {
             *sem.access_once() = sem.value;
         }
 
-        template<typename S> __device__ inline int wait(S &sem, wait_cond_t cond) {
-            int ret = 0;
-            switch(cond) {
-            case GDS_WAIT_COND_EQ:  ret = wait_eq(sem);  break;
-            case GDS_WAIT_COND_GEQ: ret = wait_geq(sem); break;
-            case GDS_WAIT_COND_AND: ret = wait_and(sem); break;
-            case GDS_WAIT_COND_NOR: ret = wait_nor(sem); break;
-            default: ret = ERROR_INVALID; break;
-            }
-            return ret;
-        }
-
         template <typename S> __device__ inline int wait_eq(S &sem) {
             int ret = ERROR_TIMEOUT;
             volatile clock_t tmout = clock() + large_timeout;
@@ -155,6 +143,17 @@ namespace gdsync {
             return ret;
         }
 
+        template<typename S> __device__ inline int wait(S &sem, wait_cond_t cond) {
+            int ret = 0;
+            switch(cond) {
+            case GDS_WAIT_COND_EQ:  ret = wait_eq(sem);  break;
+            case GDS_WAIT_COND_GEQ: ret = wait_geq(sem); break;
+            case GDS_WAIT_COND_AND: ret = wait_and(sem); break;
+            case GDS_WAIT_COND_NOR: ret = wait_nor(sem); break;
+            default: ret = ERROR_INVALID; break;
+            }
+            return ret;
+        }
     } // namespace device
 #endif
 
diff --git a/m4/ax_check_compile_flag.m4 b/m4/ax_check_compile_flag.m4
new file mode 100644
index 0000000..bd753b3
--- /dev/null
+++ b/m4/ax_check_compile_flag.m4
@@ -0,0 +1,53 @@
+# ===========================================================================
+#  https://www.gnu.org/software/autoconf-archive/ax_check_compile_flag.html
+# ===========================================================================
+#
+# SYNOPSIS
+#
+#   AX_CHECK_COMPILE_FLAG(FLAG, [ACTION-SUCCESS], [ACTION-FAILURE], [EXTRA-FLAGS], [INPUT])
+#
+# DESCRIPTION
+#
+#   Check whether the given FLAG works with the current language's compiler
+#   or gives an error.  (Warnings, however, are ignored)
+#
+#   ACTION-SUCCESS/ACTION-FAILURE are shell commands to execute on
+#   success/failure.
+#
+#   If EXTRA-FLAGS is defined, it is added to the current language's default
+#   flags (e.g. CFLAGS) when the check is done.  The check is thus made with
+#   the flags: "CFLAGS EXTRA-FLAGS FLAG".  This can for example be used to
+#   force the compiler to issue an error when a bad flag is given.
+#
+#   INPUT gives an alternative input source to AC_COMPILE_IFELSE.
+#
+#   NOTE: Implementation based on AX_CFLAGS_GCC_OPTION. Please keep this
+#   macro in sync with AX_CHECK_{PREPROC,LINK}_FLAG.
+#
+# LICENSE
+#
+#   Copyright (c) 2008 Guido U. Draheim <guidod@gmx.de>
+#   Copyright (c) 2011 Maarten Bosmans <mkbosmans@gmail.com>
+#
+#   Copying and distribution of this file, with or without modification, are
+#   permitted in any medium without royalty provided the copyright notice
+#   and this notice are preserved.  This file is offered as-is, without any
+#   warranty.
+
+#serial 6
+
+AC_DEFUN([AX_CHECK_COMPILE_FLAG],
+[AC_PREREQ(2.64)dnl for _AC_LANG_PREFIX and AS_VAR_IF
+AS_VAR_PUSHDEF([CACHEVAR],[ax_cv_check_[]_AC_LANG_ABBREV[]flags_$4_$1])dnl
+AC_CACHE_CHECK([whether _AC_LANG compiler accepts $1], CACHEVAR, [
+  ax_check_save_flags=$[]_AC_LANG_PREFIX[]FLAGS
+  _AC_LANG_PREFIX[]FLAGS="$[]_AC_LANG_PREFIX[]FLAGS $4 $1"
+  AC_COMPILE_IFELSE([m4_default([$5],[AC_LANG_PROGRAM()])],
+    [AS_VAR_SET(CACHEVAR,[yes])],
+    [AS_VAR_SET(CACHEVAR,[no])])
+  _AC_LANG_PREFIX[]FLAGS=$ax_check_save_flags])
+AS_VAR_IF(CACHEVAR,yes,
+  [m4_default([$2], :)],
+  [m4_default([$3], :)])
+AS_VAR_POPDEF([CACHEVAR])dnl
+])dnl AX_CHECK_COMPILE_FLAGS