diff --git a/runtime/README.md b/runtime/README.md
index 0328c5eb..55edff34 100644
--- a/runtime/README.md
+++ b/runtime/README.md
@@ -29,11 +29,17 @@
   * No thread-safe access (unless all required arrays are present)
   * No standard-compliant runtime API (might change long term) 
 
+  Requirements:
+  * `python3` with `jinja2` package is required to generate some of the interfaces
+     this runtime is using.
+  * Non-standard GNU Fortran extension `sizeof` must be supported by the used Fortran compiler
+  * Assumed-type, assumed-rank arrays must be supported by the used Fortran compiler
+
+
 ## Runtime subfolders
 
 * `openacc_gomp`: contains extended interfaces to GCC's LIBGOMP.
-* `gpufortrt`: contains a minimal non-standard-compliant runtime written 
-  completely in Fortran, which can be used teaching purposes
+* `gpufortrt`: contains a minimal non-standard-compliant runtime.
 
 ## Building
 
@@ -43,22 +49,6 @@ make
 ```
 
 Take a look into the `Makefile` in each folder and adapt it to your needs.
-Certain Makefile variables can be overwritten via the environment variables  
-`FC`, `HIPFORT_INC`, `FCFLAGS`, `CXX`, `CXXFLAGS`.
-
-Example:
-
-```
-cd <runtime_subfolder>
-FC=/usr/bin/gfortran HIPFORT_INC=/opt/rocm/hipfort/install/include/ make clean all
-```
-
-> **NOTE:** `python3` with `jinja2` package installed is required to use the codegen target.`
+Certain Makefile variables can be overwritten via environment variables.
 
-## Outlook
 
-* We require a standard-compliant, thread-safe HIP C++ based runtime that can
-  be used/built for HIP devices from different vendors (AMD, NVIDIA) to
-  support efforts such as GPUFORT. Main purpose of the runtime must be
-  mapping data between host and device and managing HIP streams.
-  Other features are optional.
\ No newline at end of file
diff --git a/runtime/gpufortrt/Makefile b/runtime/gpufortrt/Makefile
index 9242febb..eb647625 100644
--- a/runtime/gpufortrt/Makefile
+++ b/runtime/gpufortrt/Makefile
@@ -2,7 +2,7 @@
 # Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
 include rules.mk
 
-CXXFLAGS += -I./src -I./src/internal
+CXXFLAGS += -I./src -I./src/internal -I./include
 
 VPATH = ./src:./src/internal
 
@@ -36,7 +36,7 @@ codegen:
 
 clean_all:
 	rm -f *.o *.mod *.a
-	rm -rf include/ lib/
+	rm -rf include/*.mod lib/
 
 clean:
 	rm -f *.o *.mod *.a
diff --git a/runtime/gpufortrt/src/gpufortrt_api.h b/runtime/gpufortrt/include/gpufortrt_api.h
similarity index 85%
rename from runtime/gpufortrt/src/gpufortrt_api.h
rename to runtime/gpufortrt/include/gpufortrt_api.h
index 3cecc519..81bd6d32 100644
--- a/runtime/gpufortrt/src/gpufortrt_api.h
+++ b/runtime/gpufortrt/include/gpufortrt_api.h
@@ -5,10 +5,22 @@
 #ifdef __cplusplus
 extern "C" {
 #endif
-  extern int gpufortrt_async_noval;
-
+  int gpufortrt_get_num_devices();
+  
+  void gpufortrt_set_device_num(int dev_num);
+  int gpufortrt_get_device_num();
+  
+  size_t gpufortrt_get_property(int dev_num,
+                                gpufortrt_device_property_t property);
+  const 
+  char* gpufortrt_get_property_string(int dev_num,
+                                      gpufortrt_device_property_t property);
+  
   void gpufortrt_init();
   void gpufortrt_shutdown();
+
+  int gpufortrt_get_default_async(void);
+  void gpufortrt_set_default_async(int async_arg);
   
   void gpufortrt_data_start(
           gpufortrt_mapping_t* mappings,
@@ -77,7 +89,7 @@ extern "C" {
           void* hostptr,
           std::size_t num_bytes,
           bool never_deallocate);
-  void* gpufortrt_create_async(
+  void gpufortrt_create_async(
           void* hostptr,
           std::size_t num_bytes,
           int async_arg,
@@ -87,21 +99,11 @@ extern "C" {
           void* hostptr,
           std::size_t num_bytes,
           bool never_deallocate);
-  void* gpufortrt_copyin_async(
+  void gpufortrt_copyin_async(
           void* hostptr,
           std::size_t num_bytes,
           int async_arg,
           bool never_deallocate);
-  
-  void* gpufortrt_copy(
-         void* hostptr,
-         std::size_t num_bytes,
-         bool never_deallocate);
-  void* gpufortrt_copy_async(
-         void* hostptr,
-         std::size_t num_bytes,
-         int async_arg,
-         bool never_deallocate);
 
   // other runtime calls
 
@@ -138,6 +140,11 @@ extern "C" {
                             bool if_arg);
   void gpufortrt_wait_all_async(int* async_arg,int num_async,
                                 bool if_arg);
+  
+  int gpufortrt_async_test(int wait_arg);
+  int gpufortrt_async_test_device(int wait_arg, int dev_num);
+  int gpufortrt_async_test_all(void);
+  int gpufortrt_async_test_all_device(int dev_num);
 
   gpufortrt_queue_t gpufortrt_get_stream(int async_arg);
  
diff --git a/runtime/gpufortrt/src/gpufortrt_types.h b/runtime/gpufortrt/include/gpufortrt_types.h
similarity index 73%
rename from runtime/gpufortrt/src/gpufortrt_types.h
rename to runtime/gpufortrt/include/gpufortrt_types.h
index a52b397c..d05b17e0 100644
--- a/runtime/gpufortrt/src/gpufortrt_types.h
+++ b/runtime/gpufortrt/include/gpufortrt_types.h
@@ -8,6 +8,19 @@
 #ifdef __cplusplus
 extern "C" {
 #endif
+  extern int gpufortrt_async_noval;
+  extern int gpufortrt_async_sync;
+  extern int gpufortrt_async_default;
+  
+  enum gpufortrt_property_t {
+    gpufortrt_property_memory = 0,//>integer,  size of device memory in bytes
+    gpufortrt_property_free_memory,//>integer,  free device memory in bytes
+    gpufortrt_property_shared_memory_support,//>integer,  nonzero if the specified device supports sharing memory with the local thread
+    gpufortrt_property_name,//>string, device name*/
+    gpufortrt_property_vendor,//>string, device vendor*/
+    gpufortrt_property_driver,//>string, device driver version*/
+  };
+
   /**
    * \note: Enum values must match those of Fortran enumeration!
    * \note: Upper case first letter used because `delete` is C++ keyword.
@@ -56,4 +69,4 @@ std::ostream& operator<<(std::ostream& os, gpufortrt_map_kind_t map_kind);
 std::ostream& operator<<(std::ostream& os, gpufortrt_counter_t counter);
 std::ostream& operator<<(std::ostream& os,const gpufortrt_mapping_t& mapping);
 #endif
-#endif // GPUFORTRT_TYPES_H
\ No newline at end of file
+#endif // GPUFORTRT_TYPES_H
diff --git a/runtime/gpufortrt/include/openacc.h b/runtime/gpufortrt/include/openacc.h
new file mode 100644
index 00000000..ac6ab14d
--- /dev/null
+++ b/runtime/gpufortrt/include/openacc.h
@@ -0,0 +1,167 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
+#ifndef OPENACC_LIB_H
+#define OPENACC_LIB_H
+#ifdef __cplusplus
+extern "C" {
+#endif
+#include "gpufortrt_types.h"
+
+extern int acc_async_noval;
+extern int acc_async_sync;
+extern int acc_async_default;
+
+#define h_void void
+#define d_void void
+
+/** \note Enum values assigned according to `acc_set_device_num` description.*/
+enum acc_device_t {
+  acc_device_default = -1,
+  acc_device_all = 0,     
+  acc_device_none = 1,
+  acc_device_host,
+  acc_device_current,
+  acc_device_not_host,
+  acc_device_hip = acc_device_not_host,
+  acc_device_radeon = acc_device_hip,
+  acc_device_nvidia = acc_device_hip
+};
+
+enum acc_property_t {
+  acc_property_memory = 0,//>integer,  size of device memory in bytes
+  acc_property_free_memory,//>integer,  free device memory in bytes
+  acc_property_shared_memory_support,//>integer,  nonzero if the specified device supports sharing memory with the local thread
+  acc_property_name,//>string, device name*/
+  acc_property_vendor,//>string, device vendor*/
+  acc_property_driver//>string, device driver version*/
+};
+
+const char* ACC_DEVICE_TYPE_ENV_VAR = "ACC_DEVICE_TYPE"; //> Upper case identifer, maps to `acc_device_t` suffix.
+const char* ACC_DEVICE_NUM_ENV_VAR = "ACC_DEVICE_NUM"; //> Non-negative integer, number of default devices.
+//const char* ACC_PROFLIB_ENV_VAR = "ACC_PROFLIB";
+
+#define h_void void
+#define d_void void
+
+// except for acc_on_device all routines are only available on the host
+
+int acc_get_num_devices(acc_device_t dev_type);
+
+void acc_set_device_type(acc_device_t dev_type);
+acc_device_t acc_get_device_type(void);
+
+void acc_set_device_num(int dev_num, acc_device_t dev_type);
+int acc_get_device_num(acc_device_t dev_type);
+
+size_t acc_get_property(int dev_num,
+                        acc_device_t dev_type,
+                        acc_device_property_t property);
+const 
+char* acc_get_property_string(int dev_num,
+                              acc_device_t dev_type,
+                              acc_device_property_t property);
+
+void acc_init(acc_on_device_t dev_type);
+void acc_shutdown(acc_device_t dev_type);
+
+int acc_async_test(int wait_arg);
+int acc_async_test_device(int wait_arg, int dev_num);
+int acc_async_test_all(void);
+int acc_async_test_all_device(int dev_num);
+
+void acc_wait(int wait_arg);
+void acc_wait_device(int wait_arg, int dev_num);
+void acc_wait_async(int wait_arg, int async_arg);
+void acc_wait_device_async(int wait_arg, int async_arg,
+int dev_num);
+void acc_wait_all(void);
+void acc_wait_all_device(int dev_num);
+void acc_wait_all_async(int async_arg);
+void acc_wait_all_device_async(int async_arg, int dev_num);
+
+int acc_get_default_async(void);
+void acc_set_default_async(int async_arg);
+
+#ifdef __HIPCC__
+__host__ __device__ int acc_on_device(acc_device_t dev_type);
+#else
+int acc_on_device(acc_device_t dev_type);
+#endif
+
+d_void* acc_malloc(size_t bytes); // TODO create gpufortrt equivalent
+void acc_free(d_void* data_dev);  // TODO create gpufortr equivalent
+
+d_void* acc_copyin(h_void* data_arg, size_t bytes);
+void acc_copyin_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+
+d_void* acc_create(h_void* data_arg, size_t bytes);
+void acc_create_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+
+void acc_copyout(h_void* data_arg, size_t bytes);
+void acc_copyout_async(h_void* data_arg, size_t bytes,
+                       int async_arg);
+void acc_copyout_finalize(h_void* data_arg, size_t bytes);
+void acc_copyout_finalize_async(h_void* data_arg, size_t bytes,
+                                int async_arg);
+
+void acc_delete(h_void* data_arg, size_t bytes);
+void acc_delete_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+void acc_delete_finalize(h_void* data_arg, size_t bytes);
+void acc_delete_finalize_async(h_void* data_arg,
+                               size_t bytes, int async_arg);
+
+void acc_update_device(h_void* data_arg, size_t bytes);
+void acc_update_device_async(h_void* data_arg, size_t bytes,
+                             int async_arg);
+
+void acc_update_self(h_void* data_arg, size_t bytes);
+void acc_update_self_async(h_void* data_arg, size_t bytes,
+                           int async_arg);
+
+void acc_map_data(h_void* data_arg, d_void* data_dev,
+                  size_t bytes);
+void acc_unmap_data(h_void* data_arg);
+
+d_void* acc_deviceptr(h_void* data_arg);
+
+h_void* acc_hostptr(d_void* data_dev);
+
+int acc_is_present(h_void* data_arg, size_t bytes);
+
+void acc_memcpy_to_device(d_void* data_dev_dest,
+                          h_void* data_host_src, size_t bytes);
+void acc_memcpy_to_device_async(d_void* data_dev_dest,
+                                h_void* data_host_src, size_t bytes,
+                                int async_arg);
+
+void acc_memcpy_from_device(h_void* data_host_dest,
+                            d_void* data_dev_src, size_t bytes);
+void acc_memcpy_from_device_async(h_void* data_host_dest,
+                                  d_void* data_dev_src, size_t bytes,
+                                  int async_arg)
+
+void acc_attach(h_void** ptr_addr);
+void acc_attach_async(h_void** ptr_addr, int async_arg);
+
+void acc_detach(h_void** ptr_addr);
+void acc_detach_async(h_void** ptr_addr, int async_arg);
+void acc_detach_finalize(h_void** ptr_addr);
+void acc_detach_finalize_async(h_void** ptr_addr,
+                               int async_arg);
+
+void acc_memcpy_d2d(h_void* data_arg_dest,
+                    h_void* data_arg_src, size_t bytes,
+                    int dev_num_dest, int dev_num_src);
+
+void acc_memcpy_d2d_async(h_void* data_arg_dest,
+                          h_void* data_arg_src, size_t bytes,
+                          int dev_num_dest, int dev_num_src,
+                          int async_arg_src);
+
+#ifdef __cplusplus
+} // extern "C"
+#endif
+#endif // OPENACC_LIB_H
diff --git a/runtime/gpufortrt/src/gpufortrt_api.cpp b/runtime/gpufortrt/src/gpufortrt_api.cpp
index 17282ed4..c7dc4fef 100644
--- a/runtime/gpufortrt/src/gpufortrt_api.cpp
+++ b/runtime/gpufortrt/src/gpufortrt_api.cpp
@@ -9,8 +9,73 @@
 
 #include "internal/auxiliary.h"
 #include "internal/gpufortrt_core.h"
+  
+int gpufortrt_get_default_async(void) {
+  return gpufortrt::internal::default_async_arg;  
+}
+
+void gpufortrt_set_default_async(int async_arg) {
+  gpufortrt::internal::default_async_arg = async_arg;
+}
+
+void gpufortrt_set_device_num(int dev_num, gpufortrt_device_t dev_type) {
+  HIP_CHECK(hipSetDevice(dev_num)) // TODO backend specific, externalize
+}
+
+int gpufortrt_get_device_num(gpufortrt_device_t dev_type) {
+  int dev_num;
+  HIP_CHECK(hipGetDevice(&dev_num))
+  return dev_num;
+}
 
-int gpufortrt_async_noval = -1;
+size_t gpufortrt_get_property(int dev_num,
+                              gpufortrt_device_property_t property) {
+  switch ( property ) {
+    case gpufortrt_property_memory:
+      size_t free;
+      size_t total;
+      HIP_CHECK(hipMemGetInfo(&free, &total))
+      return total;
+      break;
+    case gpufortrt_free_memory:
+      size_t free;
+      size_t total;
+      HIP_CHECK(hipMemGetInfo(&free, &total))
+      return free;
+      break;
+    case gpufortrt_property_shared_memory_support:
+      int result;
+      HIP_CHECK(hipDeviceGetAttribute(&result,hipDeviceAttributeManagedMemory,dev_num))
+      return result;
+      break;
+    default:
+      throw std::invalid_argument("gpufortrt_get_property: property type must be 'gpufortrt_property_memory', 'gpufortrt_free_memory', or 'gpufortrt_property_shared_memory_support'");
+      break;
+  }
+}
+
+const 
+char* gpufortrt_get_property_string(int dev_num,
+                                    gpufortrt_device_property_t property) {
+      throw std::invalid_argument("gpufortrt_get_property_string: not implemented"); // TODO implement
+}
+
+// Explicit Fortran interfaces that assume device number starts from 1
+void gpufortrt_set_device_num_f(int dev_num, gpufortrt_device_t dev_type) {
+  gpufortrt_set_device_num(dev_num-1,dev_type);
+}
+int gpufortrt_get_device_num_f(gpufortrt_device_t dev_type) {
+  return gpufortrt_get_device_num(dev_type)+1;
+}
+size_t gpufortrt_get_property_f(int dev_num,
+                                gpufortrt_device_property_t property) {
+  return gpufortrt_get_property(dev_num-1,property);
+}
+const 
+char* gpufortrt_get_property_string_f(int dev_num,
+                                      gpufortrt_device_property_t property) {
+  return gpufortrt_get_property_string(dev_num-1,property);
+}
 
 void gpufortrt_mapping_init(
     gpufortrt_mapping_t* mapping,
@@ -60,8 +125,26 @@ void gpufortrt_shutdown() {
   gpufortrt::internal::record_list.destroy();
   gpufortrt::internal::queue_record_list.destroy();
 }
+  
+namespace gpufortrt {
+  namespace internal {
+      std::tuple<bool/*blocking*/,int/*async_arg*/> check_async_arg(const int async_arg) {
+      if ( async_arg >= gpufortrt_async_sync ) {
+        return std::make_tuple(
+                async_arg == gpufortrt_async_sync/*blocking*/,
+                async_arg);
+      } else {
+        int default_async_arg = gpufortrt::internal::default_async_arg;
+        return std::make_tuple(
+                default_async_arg == gpufortrt_async_sync/*blocking*/,
+                default_async_arg);
+      }
+    }
+  }
+}
 
 namespace {
+
   void* no_create_action(const gpufortrt_counter_t ctr_to_update,
                          void* hostptr,
                          std::size_t num_bytes) {
@@ -260,10 +343,15 @@ void gpufortrt_data_end() {
 void gpufortrt_data_start_async(gpufortrt_mapping_t* mappings,int num_mappings,int async_arg) {
   if ( !gpufortrt::internal::initialized ) gpufortrt_init();
   gpufortrt::internal::structured_region_stack.enter_structured_region();
+  
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::apply_mappings(mappings,
                    num_mappings,
                    gpufortrt_counter_structured,
-                   false,async_arg,false/*finalize*/);
+                   blocking,/*blocking*/
+                   async_val,
+                   false/*finalize*/);
 }
 
 void gpufortrt_data_end_async(int async_arg) {
@@ -288,10 +376,13 @@ void gpufortrt_enter_exit_data_async(gpufortrt_mapping_t* mappings,
                                      int async_arg,
                                      bool finalize) {
   if ( !gpufortrt::internal::initialized ) gpufortrt_init();
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::apply_mappings(mappings,
                  num_mappings,
                  gpufortrt_counter_dynamic,
-                 false,async_arg,
+                 blocking,/*blocking*/
+                 async_val,
                  finalize);
 }
 
@@ -302,7 +393,7 @@ void* gpufortrt_present(void* hostptr,std::size_t num_bytes) {
     num_bytes,
     gpufortrt_map_kind_present,
     false,/*never_deallocate*/
-    true/*blocking*/,
+    true,/*blocking*/
     gpufortrt_async_noval);
 }
 
@@ -317,8 +408,8 @@ void* gpufortrt_create(void* hostptr,std::size_t num_bytes,bool never_deallocate
     gpufortrt_async_noval);
 }
 
-void* gpufortrt_create_async(void* hostptr,std::size_t num_bytes,int async_arg,bool never_deallocate) {
-  return gpufortrt_create(hostptr,num_bytes,never_deallocate);
+void gpufortrt_create_async(void* hostptr,std::size_t num_bytes,int async_arg,bool never_deallocate) {
+  gpufortrt_create(hostptr,num_bytes,never_deallocate);
 }
 
 void gpufortrt_delete(void* hostptr,std::size_t num_bytes) {
@@ -336,26 +427,30 @@ void gpufortrt_delete_finalize(void* hostptr,std::size_t num_bytes) {
     num_bytes,
     gpufortrt_map_kind_delete,
     true/*finalize*/,
-    true/*blocking*/,
+    true,/*blocking*/
     gpufortrt_async_noval);
 }
 void gpufortrt_delete_async(void* hostptr,std::size_t num_bytes,int async_arg) {
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::decrement_release_action(
     hostptr,
     num_bytes,
     gpufortrt_map_kind_delete,
     false/*finalize*/,
-    false/*blocking*/,
-    async_arg);
+    blocking,
+    async_val);
 }
 void gpufortrt_delete_finalize_async(void* hostptr,std::size_t num_bytes,int async_arg) {
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::decrement_release_action(
     hostptr,
     num_bytes,
     gpufortrt_map_kind_delete,
     true/*finalize*/,
-    false/*blocking*/,
-    async_arg);
+    blocking,
+    async_val);
 }
 
 void gpufortrt_copyout(void* hostptr,std::size_t num_bytes) {
@@ -364,7 +459,7 @@ void gpufortrt_copyout(void* hostptr,std::size_t num_bytes) {
     num_bytes,
     gpufortrt_map_kind_copyout,
     false/*finalize*/,
-    true/*blocking*/,
+    true,/*blocking*/
     gpufortrt_async_noval);
 }
 void gpufortrt_copyout_finalize(void* hostptr,std::size_t num_bytes) {
@@ -373,26 +468,30 @@ void gpufortrt_copyout_finalize(void* hostptr,std::size_t num_bytes) {
     num_bytes,
     gpufortrt_map_kind_copyout,
     true/*finalize*/,
-    true/*blocking*/,
+    true,/*blocking*/
     gpufortrt_async_noval);
 }
 void gpufortrt_copyout_async(void* hostptr,std::size_t num_bytes,int async_arg) {
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::decrement_release_action(
     hostptr,
     num_bytes,
     gpufortrt_map_kind_copyout,
     false/*finalize*/,
-    false/*blocking*/,
-    async_arg);
+    blocking,
+    async_val);
 }
 void gpufortrt_copyout_finalize_async(void* hostptr,std::size_t num_bytes,int async_arg) {
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
   ::decrement_release_action(
     hostptr,
     num_bytes,
     gpufortrt_map_kind_copyout,
     true/*finalize*/,
-    false/*blocking*/,
-    async_arg);
+    blocking,
+    async_val);
 }
 
 void* gpufortrt_copyin(void* hostptr,std::size_t num_bytes,bool never_deallocate) {
@@ -402,39 +501,20 @@ void* gpufortrt_copyin(void* hostptr,std::size_t num_bytes,bool never_deallocate
     num_bytes,
     gpufortrt_map_kind_copyin,
     never_deallocate,
-    true/*blocking*/,
+    true,/*blocking*/
     gpufortrt_async_noval);
 }
-void* gpufortrt_copyin_async(void* hostptr,std::size_t num_bytes,int async_arg,bool never_deallocate) {
-  return ::create_increment_action(
+void gpufortrt_copyin_async(void* hostptr,std::size_t num_bytes,int async_arg,bool never_deallocate) {
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
+  create_increment_action(
     gpufortrt_counter_dynamic,
     hostptr,
     num_bytes,
     gpufortrt_map_kind_copyin,
     never_deallocate,
-    false/*blocking*/,
-    async_arg);
-}
-
-void* gpufortrt_copy(void* hostptr,std::size_t num_bytes,bool never_deallocate) {
-  return ::create_increment_action(
-    gpufortrt_counter_dynamic,
-    hostptr,
-    num_bytes,
-    gpufortrt_map_kind_copy,
-    never_deallocate,
-    true/*blocking*/,
-    gpufortrt_async_noval);
-}
-void* gpufortrt_copy_async(void* hostptr,std::size_t num_bytes,int async_arg,bool never_deallocate) {
-  return ::create_increment_action(
-    gpufortrt_counter_dynamic,
-    hostptr,
-    num_bytes,
-    gpufortrt_map_kind_copy,
-    never_deallocate,
-    false/*blocking*/,
-    async_arg);
+    blocking,
+    async_val);
 }
 
 namespace {
@@ -555,6 +635,31 @@ void gpufortrt_wait_all_async(int* async_arg,int num_async,
   }
 }
   
+int gpufortrt_async_test(int wait_arg) {
+  gpufortrt::internal::queue_record_list.synchronize(wait_arg);
+}
+
+int gpufortrt_async_test_device(int wait_arg, int dev_num) {
+  const int current_device_num = gpufortrt_get_device_num();
+  gpufortrt_set_device_num(dev_num);
+  gpufortrt_async_test(wait_arg);
+  gpufortrt_set_device_num(current_device_num);
+}
+
+int gpufortrt_async_test_all() {
+  for (size_t i = 0; i < queue_record_list.size(); i++) {
+    auto& queue = gpufortrt::internal::queue_record_list[i].queue;
+    HIP_CHECK(hipStreamQuery(queue))// TODO backend specific, externalize
+  } 
+}
+
+int gpufortrt_async_test_all_device(int dev_num) {
+  gpufortrt_set_device_num(dev_num);
+  int result = gpufortrt_async_test(dev_num);
+  gpufortrt_set_device_num(current_device_num);
+  return result;
+}
+  
 void* gpufortrt_deviceptr(void* hostptr) {
   LOG_INFO(1,"deviceptr; "
           << "; hostptr: "<<hostptr )
diff --git a/runtime/gpufortrt/src/gpufortrt_types.cpp b/runtime/gpufortrt/src/gpufortrt_types.cpp
index f9e249b5..da27d8b9 100644
--- a/runtime/gpufortrt/src/gpufortrt_types.cpp
+++ b/runtime/gpufortrt/src/gpufortrt_types.cpp
@@ -1,4 +1,7 @@
 #include "gpufortrt_types.h"
+int gpufortrt_async_noval = -1;
+int gpufortrt_async_sync = -2;
+int gpufortrt_async_default = -3;
 
 gpufortrt_queue_t gpufortrt_default_queue = nullptr;
 
@@ -35,4 +38,4 @@ std::ostream& operator<<(std::ostream& os,const gpufortrt_mapping_t& mapping) {
      << ", never_deallocate:" << mapping.never_deallocate;
   return os;
 }
-#endif
+#endif
\ No newline at end of file
diff --git a/runtime/gpufortrt/src/gpufortrt_types.f90 b/runtime/gpufortrt/src/gpufortrt_types.f90
index 0bc60d33..1c982978 100644
--- a/runtime/gpufortrt/src/gpufortrt_types.f90
+++ b/runtime/gpufortrt/src/gpufortrt_types.f90
@@ -8,23 +8,25 @@ module gpufortrt_types
   integer, parameter :: gpufortrt_handle_kind = c_int
 
   integer(gpufortrt_handle_kind), parameter :: gpufortrt_async_noval = -1 
+  integer(gpufortrt_handle_kind), parameter :: gpufortrt_async_sync = -2 
+  integer(gpufortrt_handle_kind), parameter :: gpufortrt_async_default = -3 
   
   !> Mapping kinds.
   enum, bind(c)
-    enumerator :: gpufortrt_map_kind_undefined       = 0
-    enumerator :: gpufortrt_map_kind_present         = 1
-    enumerator :: gpufortrt_map_kind_delete          = 2
-    enumerator :: gpufortrt_map_kind_create          = 3
-    enumerator :: gpufortrt_map_kind_no_create       = 4
-    enumerator :: gpufortrt_map_kind_copyin          = 5
-    enumerator :: gpufortrt_map_kind_copyout         = 6
-    enumerator :: gpufortrt_map_kind_copy            = 7
+    enumerator :: gpufortrt_map_kind_undefined = 0
+    enumerator :: gpufortrt_map_kind_present
+    enumerator :: gpufortrt_map_kind_delete
+    enumerator :: gpufortrt_map_kind_create
+    enumerator :: gpufortrt_map_kind_no_create
+    enumerator :: gpufortrt_map_kind_copyin
+    enumerator :: gpufortrt_map_kind_copyout
+    enumerator :: gpufortrt_map_kind_copy
   end enum
   
   enum, bind(c)
-    enumerator :: gpufortrt_counter_none       = 0
-    enumerator :: gpufortrt_counter_structured = 1
-    enumerator :: gpufortrt_counter_dynamic    = 2
+    enumerator :: gpufortrt_counter_none = 0
+    enumerator :: gpufortrt_counter_structured
+    enumerator :: gpufortrt_counter_dynamic
   end enum
   
   type, bind(c) :: gpufortrt_mapping_t
@@ -57,4 +59,4 @@ subroutine gpufortrt_mapping_init(mapping,hostptr,num_bytes,map_kind,never_deall
     mapping%never_deallocate = opt_never_deallocate
   end subroutine
 
-end module
\ No newline at end of file
+end module
diff --git a/runtime/gpufortrt/src/internal/gpufortrt_core.cpp b/runtime/gpufortrt/src/internal/gpufortrt_core.cpp
index fc096157..760be6e1 100644
--- a/runtime/gpufortrt/src/internal/gpufortrt_core.cpp
+++ b/runtime/gpufortrt/src/internal/gpufortrt_core.cpp
@@ -16,10 +16,16 @@ double gpufortrt::internal::REUSE_THRESHOLD = 0.9;
 int gpufortrt::internal::NUM_REFS_TO_DEALLOCATE = -5;
 
 // global variables
+int gpufortrt::internal::default_async_arg = gpufortrt_async_noval;
+
 bool gpufortrt::internal::initialized = false;
+
 std::size_t gpufortrt::internal::num_records = 0; 
+
 gpufortrt::internal::record_list_t gpufortrt::internal::record_list;
+
 gpufortrt::internal::queue_record_list_t gpufortrt::internal::queue_record_list;
+
 gpufortrt::internal::structured_region_stack_t gpufortrt::internal::structured_region_stack;
 
 bool gpufortrt::internal::implies_allocate_device_buffer(
diff --git a/runtime/gpufortrt/src/internal/gpufortrt_core.h b/runtime/gpufortrt/src/internal/gpufortrt_core.h
index c5acd9de..53eaec78 100644
--- a/runtime/gpufortrt/src/internal/gpufortrt_core.h
+++ b/runtime/gpufortrt/src/internal/gpufortrt_core.h
@@ -8,7 +8,7 @@
 #include <tuple>
 #include <cstddef>
 
-#include "../gpufortrt_types.h"
+#include "gpufortrt_types.h"
 
 #include "auxiliary.h"
 
@@ -331,8 +331,15 @@ namespace gpufortrt {
       /** Destroy this queue record,
        * set the `id` to negative value. */
       void destroy();
-      /** If this queue is initialized or destroyed. */
+      /** \return If this queue is initialized or destroyed. */
       bool is_initialized() const;
+
+      /** \return If all operations in this queue
+       * have been completed.*/
+      bool test();
+
+      /** Synchronize the queue. */
+      void synchronize();
     };
 
     struct queue_record_list_t {
@@ -358,6 +365,20 @@ namespace gpufortrt {
        * \return a queue if `id` is greater than 0, or `nullptr`.
        */
       gpufortrt_queue_t use_create_queue(const int id);
+      
+      /** \return If all operations in the queue with identifier `id` 
+       * have been completed.*/
+      bool test(const int id);
+
+      /** Synchronize queue with identifier `id`. */
+      void synchronize(const int id);
+      
+      /** \return If all operations in all queues
+       * have been completed.*/
+      bool test_all();
+
+      /** Synchronize all queues. */
+      void synchronize_all();
     };
 
     // global parameters, influenced by environment variables
@@ -371,6 +392,7 @@ namespace gpufortrt {
 
     // global variables
     extern bool initialized;
+    extern int default_async_arg;
     extern std::size_t num_records;
     extern record_list_t record_list;
     extern queue_record_list_t queue_record_list; 
diff --git a/runtime/gpufortrt/src/internal/queue_record_list_t.cpp b/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
index 6d9ecfaa..8bce782a 100644
--- a/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
+++ b/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
@@ -47,9 +47,9 @@ std::tuple<bool,std::size_t> gpufortrt::internal::queue_record_list_t::find_reco
 
 gpufortrt_queue_t gpufortrt::internal::queue_record_list_t::use_create_queue(int id) {
   if ( id > 0 ) { 
-    auto queues_tuple/*success,loc*/ = this->find_record(id);
-    const bool& success = std::get<0>(queues_tuple);
-    const std::size_t& loc = std::get<1>(queues_tuple);
+    bool success;
+    std::size_t loc;
+    std::tie(success,loc) = this->find_record(id);
     if ( success ) {
       LOG_INFO(3,"use existing queue; " << this->records[loc])
       return this->records[loc].queue; 
@@ -65,3 +65,81 @@ gpufortrt_queue_t gpufortrt::internal::queue_record_list_t::use_create_queue(int
     return gpufortrt_default_queue;
   }
 }
+
+
+namespace {
+  void synchronize_default_queue() {
+    gpufortrt::internal::queue_record_t default_queue_record;
+    default_queue_record.setup(-1,gpufortrt_default_queue);
+    default_queue_record.synchronize();
+  }
+  
+  bool test_default_queue() {
+    gpufortrt::internal::queue_record_t default_queue_record;
+    default_queue_record.setup(-1,gpufortrt_default_queue);
+    return default_queue_record.test();
+  }
+}
+
+/** \return If all operations in the queue with identifier `id` 
+ * have been completed.*/
+bool gpufortrt::internal::queue_record_list_t::test(const int id) {
+  if ( id > 0 ) {
+    bool success;
+    std::size_t loc;
+    std::tie(success,loc) = this->find_record(id);
+    if ( success ) {
+      bool result = this->records[loc].test(); 
+      LOG_INFO(3,"<test queue; " << this->records[loc] << "; result:"<<result)
+      return result;
+    } else {
+      LOG_INFO(3,"<test queue; no queue found for id="<<id<<"; result:1")
+      return true;
+    }
+  } else {
+    bool result = ::test_default_queue();
+    LOG_INFO(3,"<test default queue; result:"<< result)
+    return result;
+  }
+}
+
+/** Synchronize queue with identifier `id`. */
+void gpufortrt::internal::queue_record_list_t::synchronize(const int id) {
+  if ( id > 0 ) {
+    bool success;
+    std::size_t loc;
+    std::tie(success,loc) = this->find_record(id);
+    if ( success ) {
+      result = this->records[loc].synchronize(); 
+      LOG_INFO(3,"synchronize queue; " << this->records[loc])
+    } else {
+      LOG_INFO(3,"synchronize queue; no queue found for id="<<id)
+    }
+  } else {
+    LOG_INFO(3,"synchronize default queue; id:" << id)
+    ::synchronize_default_queue();
+  }
+}
+
+/** \return If all operations in all queues
+ * have been completed.*/
+bool gpufortrt::internal::queue_record_list_t::test_all() {
+  bool result = true;
+  for (size_t i = 0; i < queue_record_list.size(); i++) {
+    auto& queue_record = gpufortrt::internal::queue_record_list[i];
+    result &= queue_record.test();
+  } 
+  result &= ::test_default_queue();
+  LOG_INFO(3,"<test all queues; result:" << result)
+  return result;
+}
+
+/** Synchronize all queues. */
+void gpufortrt::internal::queue_record_list_t::synchronize_all() {
+  LOG_INFO(3,"synchronize all queues;")
+  for (size_t i = 0; i < queue_record_list.size(); i++) {
+    auto& queue_record = gpufortrt::internal::queue_record_list[i];
+    queue_record.synchronize();
+  } 
+  ::synchronize_default_queue();
+}
diff --git a/runtime/gpufortrt/src/internal/queue_record_t.cpp b/runtime/gpufortrt/src/internal/queue_record_t.cpp
index 1b166028..0569e323 100644
--- a/runtime/gpufortrt/src/internal/queue_record_t.cpp
+++ b/runtime/gpufortrt/src/internal/queue_record_t.cpp
@@ -30,13 +30,24 @@ void gpufortrt::internal::queue_record_t::setup(const int id) {
   assert(id > 0);
   assert(!this->is_initialized());
   this->id = id;
-  HIP_CHECK(hipStreamCreate(&this->queue))
+  HIP_CHECK(hipStreamCreate(&this->queue)) // TODO backend-specific, externalize
   LOG_INFO(4,"create queue; " << *this)
 }
 
 void gpufortrt::internal::queue_record_t::destroy() {
   LOG_INFO(4,"destroy queue; " << *this)
   assert(this->is_initialized());
-  HIP_CHECK(hipStreamDestroy(this->queue))
+  HIP_CHECK(hipStreamDestroy(this->queue)) // TODO backend-specific, externalize
   this->id = -1;
 }
+
+bool gpufortrt::internal::queue_record_t::test() {
+  bool result = hipStreamQuery(this->queue) == hipSuccess;
+  LOG_INFO(4,"<test queue; " << *this" <<; result:" << result << )
+  return result;
+}
+
+void gpufortrt::internal::queue_record_t::synchronize() {
+  LOG_INFO(4,"synchronize queue; " << *this)
+  HIP_CHECK(hipStreamSynchronize(this->queue))
+}
diff --git a/runtime/gpufortrt/src/openacc.cpp b/runtime/gpufortrt/src/openacc.cpp
new file mode 100644
index 00000000..54114188
--- /dev/null
+++ b/runtime/gpufortrt/src/openacc.cpp
@@ -0,0 +1,115 @@
+#include "openacc.h"
+#include "gpufortrt_api.h"
+#include "auxiliary.h"
+
+int acc_async_noval = -1;
+int acc_async_sync = -2;
+int acc_async_default = -3;
+
+namespace {
+  acc_device_t default_device = acc_device_hip;
+
+  acc_device_t current_device = acc_device_none;
+
+  bool check_device_type(acc_device_t dev_type,bool allow_none=true) {
+    switch (dev_type) {
+      case device_default:
+        return ::check_device_type(::default_device);
+      case acc_device_current:
+        return ::check_device_type(::current_device);
+      case acc_device_host:
+        return acc_device_host;
+      case acc_device_not_host:
+      case acc_device_hip:
+      case acc_device_radeon:
+      case acc_device_nvidia:
+        return acc_device_hip;
+      case acc_device_none:
+        if ( !allow_none ) {
+          throw std::invalid_argument("'acc_device_none' not allowed");
+        }
+        return acc_device_none;
+      default:
+        throw std::invalid_argument("unexpected value for 'dev_type'");
+        break;
+    }
+  }
+} // namespace {
+
+void acc_set_device_type(acc_device_t dev_type) {
+  ::current_device = dev_type;
+}
+
+acc_device_t acc_get_device_type(void) {
+  return ::current_device;
+}
+
+int acc_get_num_devices(acc_device_t dev_type) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    return gpufortrt_get_device_num();
+  } else { // host
+    return 1;
+  }
+}
+
+
+void acc_set_device_num(int dev_num, acc_device_t dev_type) {
+  // OpenACC 3.1, Section 3.2.4. Description:
+  // "[...]If the value of dev_num is negative, the runtime will revert to its default behavior,
+  // which is implementation-defined. If the value of the dev_type is zero, 
+  // the selected device number will be used for all device types.
+  // Calling acc_set_device_num implies a call to acc_set_device_type(dev_type)"
+  if ( static_cast<int>(dev_type) <= 0 || 
+       ::check_device_type(dev_type) == acc_device_hip ) {
+    gpufortrt_set_device_num(dev_num);
+  }
+}
+
+int acc_get_device_num(acc_device_t dev_type) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    return gpufortrt_get_device_num();
+  } else { // host
+    return 0;
+  }
+}
+
+size_t acc_get_property(int dev_num,
+                        acc_device_t dev_type,
+                        acc_device_property_t property) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    return gpufortrt_get_property(dev_num,property);
+  } else {
+    throw std::invalid_argument("acc_get_property: only implemented for non-host device types");
+  }
+}
+const 
+char* acc_get_property_string(int dev_num,
+                              acc_device_t dev_type,
+                              acc_device_property_t property) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    return gpufortrt_get_property_string(dev_num,property);
+  } else {
+    throw std::invalid_argument("acc_get_property_string: only implemented for non-host device types");
+  }
+}
+
+// Explicit Fortran interfaces that assume device number starts from 1
+void acc_set_device_num_f(int dev_num, acc_device_t dev_type) {
+  acc_set_device_num(dev_num-1,dev_type);
+}
+int acc_get_device_num_f(acc_device_t dev_type) {
+  return acc_get_device_num(dev_type)+1;
+}
+
+
+void acc_init(acc_on_device_t dev_type) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    gpufortrt_init();
+  }
+}
+
+void acc_shutdown(acc_device_t dev_type) {
+  if ( ::check_device_type(dev_type) == acc_device_hip ) {
+    gpufortrt_shutdown();
+  }
+}
diff --git a/runtime/gpufortrt/src/openacc.f90 b/runtime/gpufortrt/src/openacc.f90
new file mode 100644
index 00000000..59c507a6
--- /dev/null
+++ b/runtime/gpufortrt/src/openacc.f90
@@ -0,0 +1,661 @@
+! SPDX-License-Identifier: MIT
+! Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
+module openacc
+  use iso_c_binding, only: c_int
+  private :: c_int
+  
+  integer, parameter :: acc_handle_kind = c_int
+
+  integer(acc_handle_kind), parameter :: acc_async_noval = -1 
+  integer(acc_handle_kind), parameter :: acc_async_sync = -2 
+  integer(acc_handle_kind), parameter :: acc_async_default = -3 
+
+  integer, parameter :: acc_device_kind = kind(acc_device_none)
+  enum, bind(c)
+    enumerator :: acc_device_none = 0
+    enumerator :: acc_device_default
+    enumerator :: acc_device_host
+    enumerator :: acc_device_not_host
+    enumerator :: acc_device_current
+    enumerator :: acc_device_hip = acc_device_not_host
+    enumerator :: acc_device_radeon = acc_device_hip
+    enumerator :: acc_device_nvidia = acc_device_hip
+  end enum
+  
+  integer, parameter :: acc_device_kind = kind(acc_device_none)
+  
+  enum, bind(c)
+    enumerator :: acc_property_memory = 0 !>integer,  size of device memory in bytes
+    enumerator :: acc_property_free_memory !>integer,  free device memory in bytes
+    enumerator :: acc_property_shared_memory_support !>integer,  nonzero if the specified device supports sharing memory with the local thread
+    enumerator :: acc_property_name !>string, device name*/
+    enumerator :: acc_property_vendor !>string, device vendor*/
+    enumerator :: acc_property_driver !>string, device driver version*/
+  end enum
+
+  integer, parameter :: acc_property_kind = kind(acc_property_memory)
+
+  ! direct delegation to C implementation
+  interface
+    subroutine acc_set_device_type(dev_type) bind(c,name="acc_set_device_type")
+      implicit none
+      integer(acc_device_kind),value :: dev_type
+    end subroutine
+  
+    integer(acc_device_kind) function acc_get_device_type() bind(c,name="acc_get_device_type")
+      implicit none
+    end function
+
+    function acc_get_property(dev_num, dev_type, property) bind(c,name="acc_get_property")
+      use iso_c_binding, only: c_size_t
+      implicit none
+      integer, value :: dev_num
+      integer(acc_device_kind), value :: dev_type
+      integer(acc_device_property_kind), value :: property
+      integer(c_size_t) :: acc_get_property
+    end function
+
+    subroutine acc_init(dev_type) bind(c,name="acc_init")
+      implicit none
+      integer(acc_device_kind),value :: dev_type
+    end subroutine
+
+    subroutine acc_shutdown(dev_type) bind(c,name="acc_shutdown")
+      implicit none
+      integer(acc_device_kind),value :: dev_type
+    end subroutine
+
+    subroutine acc_wait(wait_arg) bind(c,name="acc_wait")
+      implicit none
+      integer(acc_handle_kind),value :: wait_arg
+    end subroutine
+
+    subroutine acc_wait_all() bind(c,name="acc_wait_all")
+      implicit none
+    end subroutine
+    
+    subroutine acc_wait_all_async(wait_arg) bind(c,name="acc_wait_all_async")
+      implicit none
+      integer(acc_handle_kind),value :: wait_arg
+    end subroutine
+
+    function acc_get_default_async() bind(c,name="acc_get_default_async")
+      implicit none
+      integer(acc_handle_kind),value :: acc_get_default_async
+    end function
+
+    subroutine acc_set_default_async(async_arg) bind(c,name="acc_set_default_async")
+      implicit none
+      integer(acc_handle_kind),value :: acc_get_default_async
+    end subroutine 
+
+  end interface
+
+  ! Overloaded routines 
+  
+  interface acc_copyin
+    module procedure :: acc_copyin_nb
+    module procedure :: acc_copyin_b
+  end interface
+  
+  interface acc_copyin_async
+    module procedure :: acc_copyin_async_nb
+    module procedure :: acc_copyin_async_b
+  end interface
+  
+  interface acc_create
+    module procedure :: acc_create_nb
+    module procedure :: acc_create_b
+  end interface
+  
+  interface acc_create_async
+    module procedure :: acc_create_async_nb
+    module procedure :: acc_create_async_b
+  end interface
+
+  interface acc_update_device
+    module procedure :: acc_update_device_nb
+    module procedure :: acc_update_device_b
+  end interface
+  
+  interface acc_update_device_async
+    module procedure :: acc_update_device_async_nb
+    module procedure :: acc_update_device_async_b
+  end interface
+  
+  interface acc_update_self
+    module procedure :: acc_update_self_nb
+    module procedure :: acc_update_self_b
+  end interface
+  
+  interface acc_update_self_async
+    module procedure :: acc_update_self_async_nb
+    module procedure :: acc_update_self_async_b
+  end interface
+  
+  interface acc_is_present
+    module procedure :: acc_is_present_nb
+    module procedure :: acc_is_present_b
+  end interface
+
+  ! Legacy interfaces
+  interface acc_present_or_copyin
+    module procedure :: acc_copyin_nb
+    module procedure :: acc_copyin_b
+  end interface
+  interface acc_pcopyin
+    module procedure :: acc_copyin_nb
+    module procedure :: acc_copyin_b
+  end interface
+  
+  interface acc_present_or_create
+    module procedure :: acc_create_nb
+    module procedure :: acc_create_b
+  end interface
+  interface acc_pcreate
+    module procedure :: acc_create_nb
+    module procedure :: acc_create_b
+  end interface
+
+  ! Non-standard extensions
+
+  interface acc_malloc
+    type(c_ptr) function acc_malloc_c_impl(bytes) bind(c,name="acc_malloc")
+      use iso_c_binding, only: c_size_t, c_ptr
+      implicit none
+      integer(c_size_t),value :: bytes
+    end function
+  end interface
+  
+  interface acc_free
+    subroutine acc_free_c_impl(data_dev) bind(c,name="acc_free")
+      use iso_c_binding, only: c_ptr
+      implicit none
+      type(c_ptr),value :: data_dev
+    end subroutine
+  end interface
+
+contains
+  integer function acc_get_num_devices(dev_type)
+    use iso_c_binding
+    implicit none
+    interface
+      integer(c_int) function acc_get_num_devices_c_impl(dev_type) bind(c,name="acc_get_num_devices")
+        implicit none
+        integer(acc_device_kind),value :: dev_type
+      end function
+    end interface
+    acc_get_num_devices = int(acc_get_num_devices_c_impl(dev_type))
+  end function
+    
+  subroutine acc_set_device_num(dev_num, dev_type)
+    implicit none
+    integer :: dev_num
+    integer(acc_device_kind) :: dev_type
+    interface   
+      subroutine acc_set_device_num_c_impl(dev_num, dev_type) bind(c,name="acc_set_device_num")
+        implicit none
+         integer(c_int),value :: dev_num
+         integer(acc_device_kind),value :: dev_type
+      end subroutine         
+    end interface
+    call acc_set_device_num_c_impl(int(dev_num,kind=c_int), dev_type)
+  end subroutine         
+
+  integer function acc_get_device_num(dev_type)
+    implicit none
+    integer(acc_device_kind) :: dev_type
+    interface
+      integer(c_int) function acc_get_device_num_c_impl(dev_type) bind(c,name="acc_get_device_num")
+        implicit none
+      end function
+    end interface
+    acc_get_device_num = int(acc_get_device_num_c_impl(dev_type))
+  end function
+
+  subroutine acc_get_property_string(dev_num, dev_type,&
+    implicit none
+                                     property, string)
+    integer, value :: dev_num
+    integer(acc_device_kind), value :: dev_type
+    integer(acc_device_property_kind), value :: property
+    character*(*) :: string
+    ! TODO 
+    !interface
+    !  subroutine acc_get_property_string_c_impl(dev_num,dev_type,property,string) bind(c,name="acc_get_property_string")
+    !    use iso_c_binding, only: c_int
+    !    integer(c_int), value :: dev_num
+    !    integer(acc_device_kind), value :: dev_type
+    !    integer(acc_device_property_kind), value :: property
+    !    character*(kind=c_char,len=*) :: string
+    !  end subroutine
+    !end interface
+    !character(kind=c_char,len=*) :: 
+    !call acc_get_property_string_c_impl(dev_num,dev_type,property,string)
+    ERROR STOP "acc_get_property_string not implemented yet!"
+  end subroutine
+
+  logical function acc_async_test(wait_arg)
+    implicit none
+    integer(acc_handle_kind) :: wait_arg
+    interface
+      integer(c_int) function acc_async_test_c_impl(wait_arg) bind(c,name="acc_async_test")
+        implicit none
+        integer(acc_handle_kind),value :: wait_arg
+      end function
+    end interface
+    acc_async_test = acc_async_test_c_impl(wait_arg) > 0
+  end function
+  
+  logical function acc_async_test_device(wait_arg, dev_num)
+    implicit none
+    integer(acc_handle_kind) :: wait_arg
+    integer :: dev_num
+    interface
+      integer(c_int) function acc_async_test_device_c_impl(wait_arg, dev_num) &
+              bind(c,name="acc_async_test_device")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(acc_handle_kind),value :: wait_arg
+        integer(c_int),value :: dev_num
+      end function
+    end interface
+    acc_async_test_device = acc_async_test_device_c_impl(wait_arg,int(dev_num,kind=c_int)) > 0
+  end function
+  
+  logical function acc_async_test_all()
+    implicit none
+    interface
+      integer(c_int) function acc_async_test_all_c_impl() &
+        bind(c,name=acc_async_test_all)
+        implicit none
+      end function
+    end interface
+    acc_async_test_all = acc_async_test_all_c_impl() > 0
+  end function
+  
+  logical function acc_async_test_all_device(dev_num)
+    implicit none
+    integer :: dev_num
+    interface 
+      integer(c_int) function acc_async_test_all_device_c_impl(dev_num) bind(c,name="acc_async_test_all_device")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(c_int),value :: dev_num
+      end function
+    end interface
+    acc_async_test_device = acc_async_test_all_device(int(dev_num,kind=c_int)) > 0
+  end function
+  
+  subroutine acc_wait_device(wait_arg, dev_num)
+    implicit none
+    integer(acc_handle_kind) :: wait_arg
+    integer :: dev_num
+    interface
+      subroutine acc_wait_device_c_impl(wait_arg, dev_num) &
+              bind(c,name="acc_wait_device")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(acc_handle_kind),value :: wait_arg
+        integer(c_int),value :: dev_num
+      end subroutine
+    end interface
+    call acc_wait_device_c_impl(wait_arg,int(dev_num,kind=c_int))
+  end subroutine
+  
+  subroutine acc_wait_async(wait_arg, async_arg)
+    implicit none
+    integer(acc_handle_kind) :: wait_arg, async_arg
+    interface
+      subroutine acc_wait_async_c_impl(wait_arg, async_arg) &
+              bind(c,name="acc_wait_async")
+                implicit none
+        integer(acc_handle_kind),value :: wait_arg, async_arg
+      end subroutine
+    end interface
+    call acc_wait_async_c_impl(wait_arg,async_arg)
+  end subroutine
+  
+  subroutine acc_wait_device_async(wait_arg, async_arg, dev_num)
+    implicit none
+    integer(acc_handle_kind) :: wait_arg, async_arg
+    integer :: dev_num
+    interface
+      subroutine acc_wait_device_async_c_impl(wait_arg, async_arg, dev_num) &
+              bind(c,name="acc_wait_device_async")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(acc_handle_kind),value :: wait_arg, async_arg
+        integer(c_int),value :: dev_num
+      end subroutine
+    end interface
+    call acc_wait_device_async_c_impl(wait_arg,async_arg,int(dev_num,kind=c_int))
+  end subroutine
+  
+  subroutine acc_wait_all_device(dev_num)
+    implicit none
+    integer :: dev_num
+    interface
+      subroutine acc_wait_all_device_c_impl(dev_num) &
+              bind(c,name="acc_wait_all_device")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(c_int),value :: dev_num
+      end subroutine
+    end interface
+    call acc_wait_all_device_c_impl(int(dev_num,kind=c_int))
+  end subroutine
+  
+  subroutine acc_wait_all_device_async(async_arg, dev_num)
+    implicit none
+    integer(acc_handle_kind) :: async_arg
+    integer :: dev_num
+    interface
+      subroutine acc_wait_all_device_async_c_impl(async_arg, dev_num) &
+              bind(c,name="acc_wait_all_device_async")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(acc_handle_kind),value :: async_arg
+        integer(c_int),value :: dev_num
+      end subroutine
+    end interface
+    call acc_wait_all_device_async_c_impl(async_arg,int(dev_num,kind=c_int))
+  end subroutine
+
+  logical function acc_on_device(dev_type)
+    implicit none
+    integer(acc_device_kind) :: dev_type
+    interface 
+      integer(c_int) function acc_on_device_c_impl(dev_type) bind(c,name="acc_on_device")
+        use iso_c_binding, only: c_int
+        implicit none
+        integer(acc_device_kind),value :: dev_type
+      end function
+    end interface
+    acc_on_device = acc_on_device_c_impl(dev_type) > 0
+  end function
+
+  !subroutine acc_copyin(data_arg) # TODO
+  !subroutine acc_copyin_async(data_arg, async_arg) # TODO
+  subroutine acc_copyin(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    !
+    type(c_ptr) :: tmp
+    interface
+      function acc_copyin_c_impl(data_arg, bytes) bind(c,name="acc_copyin")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        type(c_ptr) :: acc_copyin_c_impl
+      end function
+    end interface
+    tmp =  acc_copyin_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_copyin_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_copyin_async_c_impl(data_arg, bytes) bind(c,name="acc_copyin_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_copyin_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  
+  !subroutine acc_create(data_arg) # TODO
+  !subroutine acc_create_async(data_arg, async_arg) # TODO
+  subroutine acc_create(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    !
+    type(c_ptr) :: tmp
+    interface
+      function acc_create_c_impl(data_arg, bytes) bind(c,name="acc_create")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        type(c_ptr) :: acc_create_c_impl
+      end function
+    end interface
+    tmp =  acc_create_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_create_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_create_async_c_impl(data_arg, bytes) bind(c,name="acc_create_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_create_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  
+  !subroutine acc_copyout(data_arg) # TODO
+  !subroutine acc_copyout_async(data_arg, async_arg) # TODO
+  subroutine acc_copyout(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_copyout_c_impl(data_arg, bytes) bind(c,name="acc_copyout")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+      end subroutine
+    end interface
+    call acc_copyout_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_copyout_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_copyout_async_c_impl(data_arg, bytes) bind(c,name="acc_copyout_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_copyout_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  !subroutine acc_copyout_finalize(data_arg) # TODO
+  !subroutine acc_copyout_finalize_async(data_arg, async_arg) # TODO
+  subroutine acc_copyout_finalize(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_copyout_finalize_c_impl(data_arg, bytes) bind(c,name="acc_copyout_finalize")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+      end subroutine
+    end interface
+    call acc_copyout_finalize_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_copyout_finalize_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_copyout_finalize_async_c_impl(data_arg, bytes) bind(c,name="acc_copyout_finalize_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_copyout_finalize_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  
+  !subroutine acc_delete(data_arg) # TODO
+  !subroutine acc_delete_async(data_arg, async_arg) # TODO
+  subroutine acc_delete(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_delete_c_impl(data_arg, bytes) bind(c,name="acc_delete")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+      end subroutine
+    end interface
+    call acc_delete_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_delete_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_delete_async_c_impl(data_arg, bytes) bind(c,name="acc_delete_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_delete_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  !subroutine acc_delete_finalize(data_arg) # TODO
+  !subroutine acc_delete_finalize_async(data_arg, async_arg) # TODO
+  subroutine acc_delete_finalize(data_arg, bytes)
+    implicit none
+    type(*), target, dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_delete_finalize_c_impl(data_arg, bytes) bind(c,name="acc_delete_finalize")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+      end subroutine
+    end interface
+    call acc_delete_finalize_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+  subroutine acc_delete_finalize_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_delete_finalize_async_c_impl(data_arg, bytes) bind(c,name="acc_delete_finalize_async")
+        implicit none
+        type(c_ptr),value :: hostptr
+        integer(c_size_t),value :: num_bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_delete_finalize_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+  end subroutine
+
+  !subroutine acc_update_device(data_arg) # TODO
+  !subroutine acc_update_device_async(data_arg, async_arg) TODO
+  subroutine acc_update_device(data_arg, bytes)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_update_device_c_impl(data_arg, bytes) bind(c,name="acc_update_device")
+        implicit none
+        type(c_ptr),value :: data_arg
+        integer(c_size_t),value :: bytes
+      end subroutine
+    end interface
+    call acc_update_device_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
+  end subroutine
+
+  subroutine acc_update_device_async(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_update_device_c_impl(data_arg, bytes) bind(c,name="acc_update_device")
+        implicit none
+        type(c_ptr),value :: data_arg
+        integer(c_size_t),value :: bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_update_device_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t),async_arg)
+  end subroutine
+  
+  subroutine acc_update_self_b(data_arg, bytes)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    interface
+      subroutine acc_update_self_c_impl(data_arg, bytes) bind(c,name="acc_update_self")
+        implicit none
+        type(c_ptr),value :: data_arg
+        integer(c_size_t),value :: bytes
+      end subroutine
+    end interface
+    call acc_update_self_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
+  end subroutine
+
+  subroutine acc_update_self_async_b(data_arg, bytes, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    integer(acc_handle_kind) :: async_arg
+    interface
+      subroutine acc_update_self_c_impl(data_arg, bytes) bind(c,name="acc_update_self")
+        implicit none
+        type(c_ptr),value :: data_arg
+        integer(c_size_t),value :: bytes
+        integer(acc_handle_kind),value :: async_arg
+      end subroutine
+    end interface
+    call acc_update_self_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t),async_arg)
+  end subroutine
+  
+  subroutine acc_update_self_b(data_arg) 
+    implicit none
+    type(*), dimension(..), contiguous :: data_arg
+    call acc_update_self_b(data_arg,sizeof(data_arg))
+  end subroutine
+  subroutine acc_update_self_async_nb(data_arg, async_arg)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer(acc_handle_kind) :: async_arg
+    call acc_update_self_async_b(data_arg,sizeof(data_arg),async_arg)
+  end subroutine
+
+  logical function acc_is_present_b(data_arg, bytes)
+    implicit none
+    type(*), dimension(..) :: data_arg
+    integer :: bytes
+    interface 
+      integer(c_int) function acc_is_present_c_impl(data_arg,bytes) bind(c,name="acc_is_present")
+        implicit none
+        type(c_ptr),value :: data_arg
+        integer(c_size_t),value :: bytes
+      end function
+    end interface
+    acc_is_present_b = acc_is_present_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t)) > 0
+  end function
+ 
+  logical function acc_is_present_nb(data_arg)
+    implicit none
+    type(*), dimension(..), contiguous :: data_arg
+    acc_is_present_nb = acc_is_present_b(data_arg,sizeof(data_arg)) 
+  end function
+end module