diff --git a/.gitignore b/.gitignore
index 5a833ed0..61919b3e 100755
--- a/.gitignore
+++ b/.gitignore
@@ -13,3 +13,4 @@ gpufort.h
 gpufort_reductions.h
 render*.template.*
 render.py.in
+.vscode/settings.json
diff --git a/python/gpufort/scanner/parser.py b/python/gpufort/scanner/parser.py
index 223ea0ed..803eb90e 100644
--- a/python/gpufort/scanner/parser.py
+++ b/python/gpufort/scanner/parser.py
@@ -629,6 +629,7 @@ def is_end_statement_(tokens, kind):
                     current_linemap["statements"]):
                 try:
                     expand_statement_functions_(current_statement)
+                    original_statement = current_statement["body"]
                     original_statement_lower = current_statement["body"].lower()
                     util.logging.log_debug4(opts.log_prefix,"parse_file","parsing statement '{}' associated with lines [{},{}]".format(original_statement_lower.rstrip(),\
                         current_linemap["lineno"],current_linemap["lineno"]+len(current_linemap["lines"])-1))
diff --git a/python/gpufort/util/parsing/parsing.py b/python/gpufort/util/parsing/parsing.py
index 5c70f7cd..b78ec5ac 100644
--- a/python/gpufort/util/parsing/parsing.py
+++ b/python/gpufort/util/parsing/parsing.py
@@ -1942,7 +1942,6 @@ def is_declaration(tokens):
 def is_blank_line(statement):
     return not len(statement.strip())
 
-
 def is_fortran_directive(statement,modern_fortran):
     """If the statement is a directive."""
     return len(statement) > 2 and (modern_fortran and statement.lstrip()[0:2] == "!$" 
diff --git a/runtime/gpufortrt/Makefile b/runtime/gpufortrt/Makefile
index eb647625..f38dc86a 100644
--- a/runtime/gpufortrt/Makefile
+++ b/runtime/gpufortrt/Makefile
@@ -31,8 +31,8 @@ $(CXX_OBJ): %.cpp.o: %.cpp
 $(F_OBJ): %.o: %.f90
 	$(FC) -c $< $(FCFLAGS)
 
-codegen:
-	python3 codegen.py src/gpufortrt_api.template.f90 -d 7
+# codegen:
+# 	python3 codegen.py src/gpufortrt_api.template.f90 -d 7
 
 clean_all:
 	rm -f *.o *.mod *.a
diff --git a/runtime/gpufortrt/include/gpufortrt_api.h b/runtime/gpufortrt/include/gpufortrt_api.h
index 81bd6d32..703a97f5 100644
--- a/runtime/gpufortrt/include/gpufortrt_api.h
+++ b/runtime/gpufortrt/include/gpufortrt_api.h
@@ -9,13 +9,18 @@ extern "C" {
   
   void gpufortrt_set_device_num(int dev_num);
   int gpufortrt_get_device_num();
-  
+  // Explicit Fortran interfaces that assume device number starts from 1
+  void gpufortrt_set_device_num_f(int dev_num);
+  int gpufortrt_get_device_num_f();
+
   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);
-  
+  size_t gpufortrt_get_property_f(int dev_num, gpufortrt_device_property_t property);
+  const char* gpufortrt_get_property_string_f(int dev_num, gpufortrt_device_property_t property);
+
   void gpufortrt_init();
   void gpufortrt_shutdown();
 
@@ -84,6 +89,10 @@ extern "C" {
   void* gpufortrt_present(
           void* hostptr,
           std::size_t num_bytes);
+
+  bool gpufortrt_is_present(
+          void* hostptr,
+          std::size_t num_bytes);
   
   void* gpufortrt_create(
           void* hostptr,
@@ -140,13 +149,24 @@ extern "C" {
                             bool if_arg);
   void gpufortrt_wait_all_async(int* async_arg,int num_async,
                                 bool if_arg);
-  
+  void gpufortrt_wait_device(int* wait_arg, int num_wait, 
+                             int dev_num, bool if_arg);
+  void gpufortrt_wait_device_async(int* wait_arg, int num_wait, 
+                                   int* async_arg, int num_async, 
+                                   int dev_num, bool if_arg);
+  void gpufortrt_wait_all_device(int dev_num, bool if_arg);
+  void gpufortrt_wait_all_device_async(int* async_arg, int num_async, 
+                                 int dev_num, 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);
+  void* gpufortrt_malloc(size_t bytes);
+  void gpufortrt_free(void* data_dev);
+  void gpufortrt_map_data(void* data_arg, void* data_dev,
+                        size_t bytes);
  
   /** \return device pointer associated with `hostptr`, or nullptr.
    *  First searches through the structured region stack and then
diff --git a/runtime/gpufortrt/include/gpufortrt_types.h b/runtime/gpufortrt/include/gpufortrt_types.h
index d05b17e0..429d029b 100644
--- a/runtime/gpufortrt/include/gpufortrt_types.h
+++ b/runtime/gpufortrt/include/gpufortrt_types.h
@@ -12,7 +12,7 @@ extern "C" {
   extern int gpufortrt_async_sync;
   extern int gpufortrt_async_default;
   
-  enum gpufortrt_property_t {
+  enum gpufortrt_device_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
diff --git a/runtime/gpufortrt/include/openacc.h b/runtime/gpufortrt/include/openacc.h
index ac6ab14d..9326cda4 100644
--- a/runtime/gpufortrt/include/openacc.h
+++ b/runtime/gpufortrt/include/openacc.h
@@ -2,10 +2,11 @@
 // Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
 #ifndef OPENACC_LIB_H
 #define OPENACC_LIB_H
+#include "gpufortrt_types.h"
 #ifdef __cplusplus
 extern "C" {
 #endif
-#include "gpufortrt_types.h"
+
 
 extern int acc_async_noval;
 extern int acc_async_sync;
@@ -16,18 +17,17 @@ extern int acc_async_default;
 
 /** \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_none = 0,
+  acc_device_default,
   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
+  acc_device_current,
+  acc_device_hip,
+  acc_device_radeon,
+  acc_device_nvidia
 };
 
-enum acc_property_t {
+enum acc_device_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
@@ -51,7 +51,9 @@ 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);
+void acc_set_device_num_f(int dev_num, acc_device_t dev_type);
 int acc_get_device_num(acc_device_t dev_type);
+int acc_get_device_num_f(acc_device_t dev_type);
 
 size_t acc_get_property(int dev_num,
                         acc_device_t dev_type,
@@ -61,7 +63,14 @@ 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);
+size_t acc_get_property_f(int dev_num,
+                        acc_device_t dev_type,
+                        acc_device_property_t property);
+const 
+char* acc_get_property_string_f(int dev_num,
+                              acc_device_t dev_type,
+                              acc_device_property_t property);
+void acc_init(acc_device_t dev_type);
 void acc_shutdown(acc_device_t dev_type);
 
 int acc_async_test(int wait_arg);
@@ -72,9 +81,8 @@ 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_device_async(int wait_arg, int async_arg, int dev_num);
+void acc_wait_all();
 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);
@@ -141,7 +149,7 @@ 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)
+                                  int async_arg);
 
 void acc_attach(h_void** ptr_addr);
 void acc_attach_async(h_void** ptr_addr, int async_arg);
diff --git a/runtime/gpufortrt/openacc_library_routines.md b/runtime/gpufortrt/openacc_library_routines.md
new file mode 100644
index 00000000..c509ae13
--- /dev/null
+++ b/runtime/gpufortrt/openacc_library_routines.md
@@ -0,0 +1,57 @@
+---
+geometry: margin=2cm
+---
+
+# Implemented API
+
+| API | Lang\* | OpenACC | GPUFORTRT\*\* | Priority\*\*\* | 
+|-----|--------|---------|-------------|----------|
+|acc\_get\_num\_devices|C/C++, Fortran|implemented|implemented|high|
+|acc\_set\_device\_type|C/C++, Fortran|implemented|implemented|high|
+|acc\_get\_device\_type|C/C++, Fortran|implemented|implemented|high|
+|acc\_set\_device\_num|C/C++, Fortran|implemented|implemented||
+|acc\_get\_device\_num|C/C++, Fortran|implemented|implemented||
+|acc\_get\_property|C/C++, Fortran|implemented|implemented||
+|acc\_init|C/C++, Fortran|implemented|implemented||
+|acc\_shutdown|C/C++, Fortran|implemented|implemented||
+|acc\_async\_test|C/C++, Fortran|implemented|implemented||
+|acc\_async\_test\_device|C/C++, Fortran|implemented|implemented||
+|acc\_async\_test\_all|C/C++, Fortran|implemented|implemented||
+|acc\_async\_test\_all\_device|C/C++, Fortran|implemented|implemented||
+|acc\_wait|C/C++, Fortran|implemented|implemented||
+|acc\_wait\_device|C/C++, Fortran|implemented|implemented|high|
+|acc\_wait\_async|C/C++, Fortran|implemented|implemented||
+|acc\_wait\_device\_async|C/C++, Fortran|implemented|implemented|high|
+|acc\_wait\_all|C/C++, Fortran|implemented|implemented||
+|acc\_wait\_all\_device|C/C++, Fortran|implemented|implemented|high|
+|acc\_wait\_all\_async|C/C++, Fortran|implemented|implemented||
+|acc\_wait\_all\_device\_async|C/C++, Fortran|implemented|implemented|high|
+|acc\_get\_default\_async|C/C++, Fortran|implemented|implemented||
+|acc\_set\_default\_async|C/C++, Fortran|implemented|implemented||
+|acc\_on\_device||||low|
+|acc\_malloc||||low|
+|acc\_free||||low|
+|acc\_copyin|C/C++, Fortran|implemented|implemented||
+|acc\_create|C/C++, Fortran|implemented|implemented||
+|acc\_copyout|C/C++, Fortran|implemented|implemented||
+|acc\_delete|C/C++, Fortran|implemented|implemented||
+|acc\_update\_device|C/C++, Fortran|implemented|implemented||
+|acc\_update\_self|C/C++, Fortran|implemented|implemented||
+|acc\_map\_data||||low|
+|acc\_unmap\_data||||low|
+|acc\_deviceptr|C/C++||implemented||
+|acc\_hostptr|C/C++|||low|
+|acc\_is\_present|||implemented||
+|acc\_memcpy\_to\_device||||low|
+|acc\_memcpy\_from\_device||||low|
+|acc\_memcpy\_device||||low|
+|acc\_attach||||low|
+|acc\_detach||||low|
+|acc\_memcpy\_d2d||||low|
+
+Remarks:
+
+* \* While some APIs are exposed only to C according to the OpenACC standard, `GPUFORTRT` may expose some C interfaces also to Fortran. An \* indicates that this feature was exposed by the GPUFORTRT to Fortran despite the OpenACC standard not requiring this.
+* \*\* `GPUFORTRT` signatures are prefixd by `gpufortrt_` instead of `acc_` and the number and meaning of 
+arguments may differ compared to the OpenACC signature.
+* \*\*\* Current priorities for implementing missing APIs. This column will disappear as soon as all are implemented.
diff --git a/runtime/gpufortrt/rules.mk b/runtime/gpufortrt/rules.mk
index 43e7d9fe..26150fe6 100644
--- a/runtime/gpufortrt/rules.mk
+++ b/runtime/gpufortrt/rules.mk
@@ -2,7 +2,7 @@ HIP_PLATFORM ?= amd
 LIBGPUFORTRT  = libgpufortrt_$(HIP_PLATFORM).a
 
 FC      = gfortran -fmax-errors=5
-FCFLAGS ?= -std=f2008 -ffree-line-length-none -cpp
+FCFLAGS ?= -ffree-line-length-none -cpp
 
 #FCFLAGS += -g -ggdb -O0 -fbacktrace -fmax-errors=5 # -DDEBUG=3
 
diff --git a/runtime/gpufortrt/src/gpufortrt_api.cpp b/runtime/gpufortrt/src/gpufortrt_api.cpp
index c7dc4fef..4840cbb9 100644
--- a/runtime/gpufortrt/src/gpufortrt_api.cpp
+++ b/runtime/gpufortrt/src/gpufortrt_api.cpp
@@ -18,11 +18,17 @@ 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) {
+void gpufortrt_set_device_num(int dev_num) {
   HIP_CHECK(hipSetDevice(dev_num)) // TODO backend specific, externalize
 }
 
-int gpufortrt_get_device_num(gpufortrt_device_t dev_type) {
+int gpufortrt_get_num_devices() {
+  int dev_num;
+  HIP_CHECK(hipGetDeviceCount(&dev_num))
+  return dev_num;
+}
+
+int gpufortrt_get_device_num() {
   int dev_num;
   HIP_CHECK(hipGetDevice(&dev_num))
   return dev_num;
@@ -30,16 +36,13 @@ int gpufortrt_get_device_num(gpufortrt_device_t dev_type) {
 
 size_t gpufortrt_get_property(int dev_num,
                               gpufortrt_device_property_t property) {
+  size_t free, total;
   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;
+    case gpufortrt_property_free_memory:
       HIP_CHECK(hipMemGetInfo(&free, &total))
       return free;
       break;
@@ -60,13 +63,6 @@ char* gpufortrt_get_property_string(int dev_num,
       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);
@@ -77,6 +73,14 @@ char* gpufortrt_get_property_string_f(int dev_num,
   return gpufortrt_get_property_string(dev_num-1,property);
 }
 
+// Explicit Fortran interfaces that assume device number starts from 1
+void gpufortrt_set_device_num_f(int dev_num) {
+  gpufortrt_set_device_num(dev_num-1);
+}
+int gpufortrt_get_device_num_f() {
+  return gpufortrt_get_device_num()+1;
+}
+
 void gpufortrt_mapping_init(
     gpufortrt_mapping_t* mapping,
     void* hostptr,
@@ -409,7 +413,17 @@ void* gpufortrt_create(void* hostptr,std::size_t num_bytes,bool 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);
+  bool blocking; int async_val;
+  std::tie(blocking,async_val) = gpufortrt::internal::check_async_arg(async_arg);
+  // gpufortrt_create(hostptr,num_bytes,never_deallocate);
+  ::create_increment_action(
+    gpufortrt_counter_dynamic,
+    hostptr,
+    num_bytes,
+    gpufortrt_map_kind_create,
+    never_deallocate,/*never_deallocate*/
+    blocking,/*blocking*/
+    async_val);
 }
 
 void gpufortrt_delete(void* hostptr,std::size_t num_bytes) {
@@ -634,26 +648,61 @@ void gpufortrt_wait_all_async(int* async_arg,int num_async,
     }
   }
 }
-  
+
+void gpufortrt_wait_device(int* wait_arg, int num_wait, int dev_num, bool if_arg){
+  const int current_device_num = gpufortrt_get_device_num();
+  gpufortrt_set_device_num(dev_num);
+  gpufortrt_wait(wait_arg, num_wait, if_arg);
+  gpufortrt_set_device_num(current_device_num);
+}
+
+void gpufortrt_wait_device_async(int* wait_arg, int num_wait, 
+                                 int* async_arg, int num_async, 
+                                 int dev_num, bool if_arg){
+  const int current_device_num = gpufortrt_get_device_num();
+  gpufortrt_set_device_num(dev_num);
+  gpufortrt_wait_async(wait_arg, num_wait, async_arg, num_async, if_arg);
+  gpufortrt_set_device_num(current_device_num);
+}
+
+void gpufortrt_wait_all_device(int dev_num, bool if_arg){
+  const int current_device_num = gpufortrt_get_device_num();
+  gpufortrt_set_device_num(dev_num);
+  gpufortrt_wait_all(if_arg);
+  gpufortrt_set_device_num(current_device_num);
+}
+
+void gpufortrt_wait_all_device_async(int* async_arg, int num_async, 
+                                 int dev_num, bool if_arg){
+  const int current_device_num = gpufortrt_get_device_num();
+  gpufortrt_set_device_num(dev_num);
+  gpufortrt_wait_all_async(async_arg, num_async, if_arg);
+  gpufortrt_set_device_num(current_device_num);
+}
+
 int gpufortrt_async_test(int wait_arg) {
-  gpufortrt::internal::queue_record_list.synchronize(wait_arg);
+  return gpufortrt::internal::queue_record_list.test(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);
+  int result = gpufortrt_async_test(wait_arg);
   gpufortrt_set_device_num(current_device_num);
+  return result;
 }
 
 int gpufortrt_async_test_all() {
-  for (size_t i = 0; i < queue_record_list.size(); i++) {
+  for (size_t i = 0; i < gpufortrt::internal::queue_record_list.size(); i++) {
     auto& queue = gpufortrt::internal::queue_record_list[i].queue;
-    HIP_CHECK(hipStreamQuery(queue))// TODO backend specific, externalize
+    if(hipStreamQuery(queue) != hipSuccess) return 0;
+    // HIP_CHECK(hipStreamQuery(queue))// TODO backend specific, externalize
   } 
+  return 1;
 }
 
 int gpufortrt_async_test_all_device(int dev_num) {
+  const int current_device_num = gpufortrt_get_device_num();
   gpufortrt_set_device_num(dev_num);
   int result = gpufortrt_async_test(dev_num);
   gpufortrt_set_device_num(current_device_num);
@@ -735,3 +784,40 @@ void* gpufortrt_use_device(void* hostptr,bool if_arg,bool if_present_arg) {
     return hostptr;
   }
 }
+
+bool gpufortrt_is_present(void* hostptr,std::size_t num_bytes) {
+  if ( !gpufortrt::internal::initialized ) LOG_ERROR("gpufortrt_is_present: runtime not initialized")
+  if ( hostptr != nullptr ) { // nullptr means no-op
+    auto list_tuple/*success,loc,offset*/ = gpufortrt::internal::record_list.find_record(hostptr,num_bytes);
+    return std::get<0>(list_tuple);
+  } else{
+      return false;
+  }
+}
+
+void* gpufortrt_malloc(size_t bytes){
+  void* deviceptr = nullptr;
+  hipError_t ierr = hipMalloc(&deviceptr,bytes);      
+  if ( ierr == hipSuccess ) {
+    return deviceptr;
+  } else {
+    LOG_ERROR("gpufortrt_malloc: could not allocate memory on device")
+    return nullptr; /* terminates beforehand */
+  }
+}
+
+void gpufortrt_free(void* data_dev){
+  HIP_CHECK( hipFree(data_dev) );   
+}
+
+void gpufortrt_map_data(void* data_arg, void* data_dev,
+                  size_t bytes){
+  // TODO: Check wether data_arg is sub_section of host memory
+  if ( data_arg != nullptr ) {
+    HIP_CHECK(hipMemcpy(
+        data_arg,
+        data_dev,
+        bytes,
+        hipMemcpyDeviceToHost));
+  }
+}
\ No newline at end of file
diff --git a/runtime/gpufortrt/src/gpufortrt_api.f90 b/runtime/gpufortrt/src/gpufortrt_api.f90
new file mode 100644
index 00000000..9654e4d6
--- /dev/null
+++ b/runtime/gpufortrt/src/gpufortrt_api.f90
@@ -0,0 +1,3090 @@
+! SPDX-License-Identifier: MIT
+! Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
+module gpufortrt_api
+  use gpufortrt_types
+
+  interface
+    subroutine gpufortrt_init() bind(c,name="gpufortrt_init")
+      implicit none
+    end subroutine
+  
+    subroutine gpufortrt_shutdown() bind(c,name="gpufortrt_shutdown")
+      implicit none
+    end subroutine
+  
+    function gpufortrt_get_stream(async_arg) &
+        bind(c,name="gpufortrt_get_stream") &
+          result(stream)
+      use iso_c_binding, only: c_ptr
+      use gpufortrt_types, only: gpufortrt_handle_kind
+      implicit none
+      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      !
+      type(c_ptr) :: stream
+    end function
+  end interface
+   
+  interface gpufortrt_present
+    module procedure :: gpufortrt_present_b
+    module procedure :: gpufortrt_present_nb
+  end interface
+
+  interface gpufortrt_use_device
+    module procedure :: gpufortrt_use_device0_l1
+    module procedure :: gpufortrt_use_device0_l4
+    module procedure :: gpufortrt_use_device0_ch1
+    module procedure :: gpufortrt_use_device0_i1
+    module procedure :: gpufortrt_use_device0_i2
+    module procedure :: gpufortrt_use_device0_i4
+    module procedure :: gpufortrt_use_device0_i8
+    module procedure :: gpufortrt_use_device0_r4
+    module procedure :: gpufortrt_use_device0_r8
+    module procedure :: gpufortrt_use_device0_c4
+    module procedure :: gpufortrt_use_device0_c8
+    module procedure :: gpufortrt_use_device1_l1
+    module procedure :: gpufortrt_use_device1_l4
+    module procedure :: gpufortrt_use_device1_ch1
+    module procedure :: gpufortrt_use_device1_i1
+    module procedure :: gpufortrt_use_device1_i2
+    module procedure :: gpufortrt_use_device1_i4
+    module procedure :: gpufortrt_use_device1_i8
+    module procedure :: gpufortrt_use_device1_r4
+    module procedure :: gpufortrt_use_device1_r8
+    module procedure :: gpufortrt_use_device1_c4
+    module procedure :: gpufortrt_use_device1_c8
+    module procedure :: gpufortrt_use_device2_l1
+    module procedure :: gpufortrt_use_device2_l4
+    module procedure :: gpufortrt_use_device2_ch1
+    module procedure :: gpufortrt_use_device2_i1
+    module procedure :: gpufortrt_use_device2_i2
+    module procedure :: gpufortrt_use_device2_i4
+    module procedure :: gpufortrt_use_device2_i8
+    module procedure :: gpufortrt_use_device2_r4
+    module procedure :: gpufortrt_use_device2_r8
+    module procedure :: gpufortrt_use_device2_c4
+    module procedure :: gpufortrt_use_device2_c8
+    module procedure :: gpufortrt_use_device3_l1
+    module procedure :: gpufortrt_use_device3_l4
+    module procedure :: gpufortrt_use_device3_ch1
+    module procedure :: gpufortrt_use_device3_i1
+    module procedure :: gpufortrt_use_device3_i2
+    module procedure :: gpufortrt_use_device3_i4
+    module procedure :: gpufortrt_use_device3_i8
+    module procedure :: gpufortrt_use_device3_r4
+    module procedure :: gpufortrt_use_device3_r8
+    module procedure :: gpufortrt_use_device3_c4
+    module procedure :: gpufortrt_use_device3_c8
+    module procedure :: gpufortrt_use_device4_l1
+    module procedure :: gpufortrt_use_device4_l4
+    module procedure :: gpufortrt_use_device4_ch1
+    module procedure :: gpufortrt_use_device4_i1
+    module procedure :: gpufortrt_use_device4_i2
+    module procedure :: gpufortrt_use_device4_i4
+    module procedure :: gpufortrt_use_device4_i8
+    module procedure :: gpufortrt_use_device4_r4
+    module procedure :: gpufortrt_use_device4_r8
+    module procedure :: gpufortrt_use_device4_c4
+    module procedure :: gpufortrt_use_device4_c8
+    module procedure :: gpufortrt_use_device5_l1
+    module procedure :: gpufortrt_use_device5_l4
+    module procedure :: gpufortrt_use_device5_ch1
+    module procedure :: gpufortrt_use_device5_i1
+    module procedure :: gpufortrt_use_device5_i2
+    module procedure :: gpufortrt_use_device5_i4
+    module procedure :: gpufortrt_use_device5_i8
+    module procedure :: gpufortrt_use_device5_r4
+    module procedure :: gpufortrt_use_device5_r8
+    module procedure :: gpufortrt_use_device5_c4
+    module procedure :: gpufortrt_use_device5_c8
+    module procedure :: gpufortrt_use_device6_l1
+    module procedure :: gpufortrt_use_device6_l4
+    module procedure :: gpufortrt_use_device6_ch1
+    module procedure :: gpufortrt_use_device6_i1
+    module procedure :: gpufortrt_use_device6_i2
+    module procedure :: gpufortrt_use_device6_i4
+    module procedure :: gpufortrt_use_device6_i8
+    module procedure :: gpufortrt_use_device6_r4
+    module procedure :: gpufortrt_use_device6_r8
+    module procedure :: gpufortrt_use_device6_c4
+    module procedure :: gpufortrt_use_device6_c8
+    module procedure :: gpufortrt_use_device7_l1
+    module procedure :: gpufortrt_use_device7_l4
+    module procedure :: gpufortrt_use_device7_ch1
+    module procedure :: gpufortrt_use_device7_i1
+    module procedure :: gpufortrt_use_device7_i2
+    module procedure :: gpufortrt_use_device7_i4
+    module procedure :: gpufortrt_use_device7_i8
+    module procedure :: gpufortrt_use_device7_r4
+    module procedure :: gpufortrt_use_device7_r8
+    module procedure :: gpufortrt_use_device7_c4
+    module procedure :: gpufortrt_use_device7_c8
+  end interface
+
+contains
+  
+  type(c_ptr) function gpufortrt_deviceptr(hostptr) 
+    use iso_c_binding
+    implicit none
+    !
+    type(*), dimension(..), target, intent(in) :: hostptr
+    !
+    interface 
+    type(c_ptr) function gpufortrt_deviceptr_c_impl(hostptr) &
+      bind(c,name="gpufortrt_deviceptr") 
+      use iso_c_binding
+      implicit none
+      !
+      type(c_ptr),value,intent(in) :: hostptr
+    end function
+    end interface
+    !
+    gpufortrt_deviceptr = gpufortrt_deviceptr_c_impl(c_loc(hostptr))
+  end function
+
+  type(c_ptr) function gpufortrt_present_b(hostptr,num_bytes) 
+    use iso_c_binding
+    implicit none
+    !
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t), value,intent(in) :: num_bytes
+    !
+    interface 
+    type(c_ptr) function gpufortrt_present_b_c_impl(hostptr,num_bytes) &
+      bind(c,name="gpufortrt_present") 
+      use iso_c_binding
+      implicit none
+      !
+      type(c_ptr), value::hostptr
+      integer(c_size_t), value :: num_bytes
+    end function
+    end interface 
+    gpufortrt_present_b = gpufortrt_present_b_c_impl(c_loc(hostptr),int(num_bytes,kind=c_size_t))
+  end function
+
+  type(c_ptr) function gpufortrt_present_nb(hostptr)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..), contiguous :: hostptr
+    gpufortrt_present_nb = gpufortrt_present_b(c_loc(hostptr),int(sizeof(hostptr), kind = c_size_t))
+  end function
+
+  !> Ignore the result of a mapping routine.
+  !> \param[in] deviceptr a device pointer.
+  subroutine gpufortrt_ignore(deviceptr)
+    use iso_c_binding, only: c_ptr
+    implicit none
+    type(c_ptr),intent(in) :: deviceptr
+    ! nop
+  end subroutine
+
+  subroutine gpufortrt_wait(wait_arg,async_arg,condition)
+    use iso_c_binding
+    implicit none
+    integer(gpufortrt_handle_kind),dimension(:),target,intent(in),optional :: wait_arg,async_arg
+    logical,intent(in),optional :: condition
+    !
+    interface
+      subroutine gpufortrt_wait_all_c_impl(condition) &
+          bind(c,name="gpufortrt_wait_all")
+        use iso_c_binding
+        implicit none
+        logical(c_bool),value,intent(in):: condition 
+      end subroutine
+      subroutine gpufortrt_wait_all_async_c_impl(async_arg,num_async_args,condition) &
+          bind(c,name="gpufortrt_wait_all_async")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: async_arg
+        integer(c_int),value,intent(in) :: num_async_args
+        logical(c_bool),value,intent(in) :: condition 
+      end subroutine
+      subroutine gpufortrt_wait_c_impl(wait_arg,num_wait_args,condition) &
+          bind(c,name="gpufortrt_wait")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: wait_arg
+        integer(c_int),value,intent(in) :: num_wait_args
+        logical(c_bool),value,intent(in) :: condition 
+      end subroutine
+      subroutine gpufortrt_wait_async_c_impl(wait_arg,num_wait_args,&
+                                             async_arg,num_async_args,&
+                                             condition) &
+          bind(c,name="gpufortrt_wait_async")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: wait_arg, async_arg
+        integer(c_int),value,intent(in) :: num_wait_args, num_async_args
+        logical(c_bool),value,intent(in) :: condition 
+      end subroutine
+    end interface
+    !
+    logical(c_bool) :: opt_if_arg
+    !
+    opt_if_arg = .true._c_bool
+    if ( present(condition) ) opt_if_arg = logical(condition,kind=c_bool)
+    !
+    if ( present(wait_arg) ) then
+      if ( present(async_arg) ) then
+        call gpufortrt_wait_async_c_impl(&
+             c_loc(wait_arg),size(wait_arg,kind=c_int),&
+             c_loc(async_arg),size(async_arg,kind=c_int),&
+             opt_if_arg)
+      else
+        call gpufortrt_wait_c_impl(&
+             c_loc(wait_arg),size(wait_arg,kind=c_int),&
+             opt_if_arg)
+      endif
+    else
+      if ( present(async_arg) ) then
+        call gpufortrt_wait_all_async_c_impl(&
+             c_loc(async_arg),size(async_arg,kind=c_int),&
+             opt_if_arg)
+      else
+        call gpufortrt_wait_all_c_impl(opt_if_arg)
+      endif
+    endif
+  end subroutine
+
+  subroutine gpufortrt_data_start(mappings,async_arg)
+  !subroutine gpufortrt_data_start(device_kind,mappings,async_arg)
+    use iso_c_binding
+    implicit none
+    !integer,intent(in) :: device_kind
+    type(gpufortrt_mapping_t),dimension(:),target,intent(in),optional :: mappings
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    !
+    interface
+      subroutine gpufortrt_data_start_c_impl(mappings,num_mappings) bind(c,name="gpufortrt_data_start")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),intent(in),value :: mappings
+        integer(c_int),intent(in),value :: num_mappings
+      end subroutine
+      subroutine gpufortrt_data_start_async_c_impl(mappings,num_mappings,async_arg) bind(c,name="gpufortrt_data_start_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),intent(in),value :: mappings
+        integer(c_int),intent(in),value :: num_mappings
+        integer(gpufortrt_handle_kind),intent(in),value :: async_arg
+      end subroutine
+    end interface 
+    !
+    if ( present(async_arg) ) then
+      if ( present(mappings) ) then
+        call gpufortrt_data_start_async_c_impl(c_loc(mappings),size(mappings),&
+                                               int(async_arg,kind=c_int))
+      else
+        call gpufortrt_data_start_async_c_impl(c_null_ptr,0_c_int,&
+                                               int(async_arg,kind=c_int)) 
+      endif
+    else
+      if ( present(mappings) ) then
+        call gpufortrt_data_start_c_impl(c_loc(mappings),size(mappings))  
+      else
+        call gpufortrt_data_start_c_impl(c_null_ptr,0_c_int)  
+      endif
+    endif
+  end subroutine
+   
+  subroutine gpufortrt_data_end(async_arg)
+    implicit none
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    !
+    interface
+      subroutine gpufortrt_data_end_c_impl() bind(c,name="gpufortrt_data_end")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+      end subroutine
+      subroutine gpufortrt_data_end_async_c_impl(async_arg) bind(c,name="gpufortrt_data_end_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        integer(gpufortrt_handle_kind),intent(in),value :: async_arg
+      end subroutine
+    end interface 
+    !
+    if ( present(async_arg) ) then
+      call gpufortrt_data_end_async_c_impl(async_arg)  
+    else
+      call gpufortrt_data_end_c_impl()  
+    endif
+  end subroutine
+
+  subroutine gpufortrt_enter_exit_data(mappings,async_arg,finalize)
+    use iso_c_binding
+    implicit none
+    !
+    !integer,intent(in) :: device_kind
+    type(gpufortrt_mapping_t),dimension(:),target,intent(in),optional :: mappings
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    logical,intent(in),optional :: finalize
+    !
+    interface
+      subroutine gpufortrt_enter_exit_data_c_impl(mappings,num_mappings,finalize) bind(c,name="gpufortrt_enter_exit_data")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: mappings
+        integer(c_int),value,intent(in) :: num_mappings
+        logical(c_bool),value,intent(in) :: finalize
+      end subroutine
+      subroutine gpufortrt_enter_exit_data_async_c_impl(mappings,num_mappings,async_arg,finalize) bind(c,name="gpufortrt_enter_exit_data_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),value,intent(in) :: mappings
+        integer(c_int),value,intent(in) :: num_mappings
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+        logical(c_bool),value,intent(in) :: finalize
+      end subroutine
+    end interface 
+    !
+    logical(c_bool) :: opt_finalize
+    !
+    opt_finalize = .false._c_bool
+    if ( present(finalize) ) opt_finalize = logical(finalize,kind=c_bool)
+    !
+    if ( present(async_arg) ) then
+      if ( present(mappings) ) then
+        call gpufortrt_enter_exit_data_async_c_impl(&
+          c_loc(mappings),&
+          size(mappings,kind=c_int),&
+          async_arg,&
+          opt_finalize) 
+      else
+        call gpufortrt_enter_exit_data_async_c_impl(&
+          c_null_ptr,&
+          0_c_int,&
+          async_arg,&
+          opt_finalize) 
+      endif
+    else
+      if ( present(mappings) ) then
+        call gpufortrt_enter_exit_data_c_impl(&
+          c_loc(mappings),&
+          size(mappings,kind=c_int),&
+          opt_finalize) 
+      else
+        call gpufortrt_enter_exit_data_c_impl(&
+          c_null_ptr,&
+          0_c_int,&
+          opt_finalize) 
+      endif
+    endif
+  end subroutine
+
+  !> Lookup device pointer for given host pointer.
+  !> \param[in] condition condition that must be met, otherwise host pointer is returned. Defaults to '.true.'.
+  !> \param[in] if_present Only return device pointer if one could be found for the host pointer.
+  !>                       otherwise host pointer is returned. Defaults to '.false.'.
+  !> \note Returns a c_null_ptr if the host pointer is invalid, i.e. not C associated.
+  function gpufortrt_use_device_b(hostptr,condition,if_present) result(resultptr)
+    use iso_c_binding
+    implicit none
+    type(c_ptr),intent(in) :: hostptr
+    logical,intent(in),optional :: condition, if_present
+    !
+    type(c_ptr) :: resultptr
+    !
+    interface
+      function gpufortrt_use_device_c_impl(hostptr,condition,if_present) &
+          bind(c,name="gpufortrt_use_device") result(deviceptr)
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        logical(c_bool),value,intent(in) :: condition, if_present
+        !
+        type(c_ptr) :: deviceptr
+      end function
+    end interface
+    !
+    logical(c_bool) :: opt_if_arg, opt_if_present_arg
+    !
+    opt_if_arg = .true._c_bool
+    opt_if_present_arg = .false._c_bool
+    if ( present(condition) ) opt_if_arg = logical(condition,kind=c_bool)
+    if ( present(if_present) ) opt_if_present_arg = logical(if_present,kind=c_bool)
+    !
+    resultptr = gpufortrt_use_device_c_impl(hostptr,opt_if_arg,opt_if_present_arg)
+  end function
+
+  
+  subroutine gpufortrt_use_device0_l1(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_l1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical(c_bool),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical(c_bool),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_l4(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_l4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    logical,target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    logical,pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_ch1(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_ch1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    character(c_char),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    character(c_char),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_i1(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_i1(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int8_t),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int8_t),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_i2(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_i2(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_short),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_short),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_i4(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_i4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_int),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_int),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_i8(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_i8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    integer(c_long),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    integer(c_long),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_r4(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_r4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_float),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_float),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_r8(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_r8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    real(c_double),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    real(c_double),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_c4(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_c4(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_float_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_float_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device0_c8(resultptr,hostptr,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr
+    !
+    type(c_ptr) :: tmp_cptr
+    !
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr)
+  end subroutine
+
+  subroutine gpufortrt_use_device1_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(*)
+    integer,intent(in),optional :: sizes(1), lbounds(1)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(1), opt_lbounds(1)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device2_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,*)
+    integer,intent(in),optional :: sizes(2), lbounds(2)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(2), opt_lbounds(2)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device3_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,1:1,*)
+    integer,intent(in),optional :: sizes(3), lbounds(3)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(3), opt_lbounds(3)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device4_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(4), lbounds(4)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(4), opt_lbounds(4)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device5_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(5), lbounds(5)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(5), opt_lbounds(5)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device6_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(6), lbounds(6)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(6), opt_lbounds(6)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):)&
+        => resultptr
+  end subroutine
+
+  subroutine gpufortrt_use_device7_c8(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
+    use iso_c_binding
+    implicit none
+    complex(c_double_complex),target,intent(in) :: hostptr(1:1,1:1,1:1,1:1,1:1,1:1,*)
+    integer,intent(in),optional :: sizes(7), lbounds(7)
+    logical,intent(in),optional :: if_arg, if_present_arg
+    !
+    complex(c_double_complex),pointer,intent(inout) :: resultptr(:,:,:,:,:,:,:)
+    !
+    type(c_ptr) :: tmp_cptr
+    integer :: opt_sizes(7), opt_lbounds(7)
+    !
+    opt_sizes = 1
+    opt_lbounds = 1  
+    if ( present(sizes) ) opt_sizes = sizes
+    if ( present(lbounds) ) opt_lbounds = lbounds
+    tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
+    call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
+    resultptr(&
+      opt_lbounds(1):,&
+      opt_lbounds(2):,&
+      opt_lbounds(3):,&
+      opt_lbounds(4):,&
+      opt_lbounds(5):,&
+      opt_lbounds(6):,&
+      opt_lbounds(7):)&
+        => resultptr
+  end subroutine
+
+  logical function gpufortrt_is_present(data_arg, bytes) 
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(c_int),value,intent(in) :: bytes
+    !
+    interface
+      type(c_ptr) function gpufortrt_is_present_c_impl(data_arg, bytes) &
+        bind(c,name="gpufortrt_present") 
+          use iso_c_binding
+          implicit none
+          !
+          type(c_ptr), value::data_arg
+          integer(c_size_t), value :: bytes
+      end function
+    end interface
+    gpufortrt_is_present = c_associated(gpufortrt_is_present_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t)))
+  end function
+
+  function gpufortrt_map_present(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*),dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes 
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+   
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_present,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_no_create(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes 
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_no_create,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_create(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes 
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_create,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_copyin(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_copyin,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_copy(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_copy,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_copyout(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*),dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_copyout,opt_never_deallocate)
+  end function
+
+  function gpufortrt_map_delete(hostptr,num_bytes,never_deallocate) result(retval)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes 
+    logical,intent(in),optional :: never_deallocate
+    !
+    type(gpufortrt_mapping_t) :: retval
+    !
+    logical :: opt_never_deallocate = .false._c_bool
+    integer(c_size_t) :: opt_num_bytes
+    !
+    if( present( never_deallocate )) opt_never_deallocate = never_deallocate
+    if( present( num_bytes )) then
+      opt_num_bytes = num_bytes
+    else 
+      opt_num_bytes = int(sizeof(hostptr),kind=c_size_t)
+    endif
+
+    call gpufortrt_mapping_init(retval,c_loc(hostptr),opt_num_bytes,&
+           gpufortrt_map_kind_delete,opt_never_deallocate)
+  end function
+
+  function gpufortrt_create(hostptr,num_bytes,never_deallocate,async_arg) result(deviceptr)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in), optional :: num_bytes
+    logical, intent(in), optional :: never_deallocate
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg 
+    !
+    type(c_ptr) :: deviceptr
+    !
+    interface
+      function gpufortrt_create_c_impl(hostptr,num_bytes,never_deallocate) &
+        bind(c,name="gpufortrt_create") result(deviceptr)
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: never_deallocate
+        !
+        type(c_ptr) :: deviceptr
+      end function
+      function gpufortrt_create_async_c_impl(hostptr,num_bytes,never_deallocate,async_arg) &
+        bind(c,name="gpufortrt_create_async") result(deviceptr)
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: never_deallocate
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+        !
+        type(c_ptr) :: deviceptr
+      end function
+    end interface
+    !
+    integer(c_size_t) :: opt_num_bytes
+    logical(c_bool) :: opt_never_deallocate
+    !
+    opt_never_deallocate = .false._c_bool
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(never_deallocate) ) opt_never_deallocate = never_deallocate
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    if ( present(async_arg) ) then
+      deviceptr = gpufortrt_create_async_c_impl(c_loc(hostptr),opt_num_bytes,opt_never_deallocate,async_arg)
+    else
+      deviceptr = gpufortrt_create_c_impl(c_loc(hostptr),opt_num_bytes,opt_never_deallocate)
+    endif
+  end function
+
+  function gpufortrt_copyin(hostptr,num_bytes,never_deallocate,async_arg) result(deviceptr)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes
+    logical,intent(in),optional :: never_deallocate
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg 
+    !
+    type(c_ptr) :: deviceptr
+    !
+    interface
+      function gpufortrt_copyin_c_impl(hostptr,num_bytes,never_deallocate) &
+          bind(c,name="gpufortrt_copyin") result(deviceptr)
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: never_deallocate
+        !
+        type(c_ptr) :: deviceptr
+      end function
+      function gpufortrt_copyin_async_c_impl(hostptr,num_bytes,never_deallocate,async_arg) &
+        bind(c,name="gpufortrt_copyin_async") result(deviceptr)
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: never_deallocate
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+        !
+        type(c_ptr) :: deviceptr
+      end function
+    end interface
+    !
+    logical(c_bool) :: opt_never_deallocate
+    integer(c_size_t) :: opt_num_bytes
+    !
+    opt_never_deallocate = .false._c_bool
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(never_deallocate) ) opt_never_deallocate = never_deallocate
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    if ( present(async_arg) ) then
+      deviceptr = gpufortrt_copyin_async_c_impl(c_loc(hostptr),opt_num_bytes,opt_never_deallocate,async_arg)
+    else
+      deviceptr = gpufortrt_copyin_c_impl(c_loc(hostptr),opt_num_bytes,opt_never_deallocate)
+    endif
+  end function
+
+  subroutine gpufortrt_delete(hostptr,num_bytes,async_arg,finalize)
+    use iso_c_binding
+    use gpufortrt_types, only: gpufortrt_handle_kind
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    logical,intent(in),optional :: finalize
+    !
+    interface
+      subroutine gpufortrt_delete_c_impl(hostptr,num_bytes) &
+          bind(c,name="gpufortrt_delete")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+      end subroutine
+      subroutine gpufortrt_delete_finalize_c_impl(hostptr,num_bytes) &
+          bind(c,name="gpufortrt_delete_finalize")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+      end subroutine
+      subroutine gpufortrt_delete_async_c_impl(hostptr,num_bytes,async_arg) &
+          bind(c,name="gpufortrt_delete_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+      subroutine gpufortrt_delete_finalize_async_c_impl(hostptr,num_bytes,async_arg) &
+          bind(c,name="gpufortrt_delete_finalize_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+    end interface
+    !
+    integer(c_size_t) :: opt_num_bytes
+    !
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    if ( present(async_arg) ) then
+      if ( present(finalize) ) then
+        call gpufortrt_delete_finalize_async_c_impl(c_loc(hostptr),opt_num_bytes,async_arg)
+      else
+        call gpufortrt_delete_async_c_impl(c_loc(hostptr),opt_num_bytes,async_arg)
+      endif
+    else
+      if ( present(finalize) ) then
+        call gpufortrt_delete_finalize_c_impl(c_loc(hostptr),opt_num_bytes)
+      else
+        call gpufortrt_delete_c_impl(c_loc(hostptr),opt_num_bytes)
+      endif
+    endif
+  end subroutine
+
+  subroutine gpufortrt_copyout(hostptr,num_bytes,async_arg,finalize)
+    use iso_c_binding
+    use gpufortrt_types, only: gpufortrt_handle_kind
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    logical,intent(in),optional :: finalize
+    !
+    interface
+      subroutine gpufortrt_copyout_c_impl(hostptr,num_bytes) &
+          bind(c,name="gpufortrt_copyout")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+      end subroutine
+      subroutine gpufortrt_copyout_finalize_c_impl(hostptr,num_bytes) &
+          bind(c,name="gpufortrt_copyout_finalize")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+      end subroutine
+      subroutine gpufortrt_copyout_async_c_impl(hostptr,num_bytes,async_arg) &
+          bind(c,name="gpufortrt_copyout_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+      subroutine gpufortrt_copyout_finalize_async_c_impl(hostptr,num_bytes,async_arg) &
+          bind(c,name="gpufortrt_copyout_finalize_async")
+        use iso_c_binding
+        use gpufortrt_types, only: gpufortrt_handle_kind
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+    end interface
+    !
+    integer(c_size_t) :: opt_num_bytes
+    !
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    if ( present(async_arg) ) then
+      if ( present(finalize) ) then
+        call gpufortrt_copyout_finalize_async_c_impl(c_loc(hostptr),opt_num_bytes,async_arg)
+      else
+        call gpufortrt_copyout_async_c_impl(c_loc(hostptr),opt_num_bytes,async_arg)
+      endif
+    else
+      if ( present(finalize) ) then
+        call gpufortrt_copyout_finalize_c_impl(c_loc(hostptr),opt_num_bytes)
+      else
+        call gpufortrt_copyout_c_impl(c_loc(hostptr),opt_num_bytes)
+      endif
+    endif
+  end subroutine
+
+  subroutine gpufortrt_update_self(hostptr,num_bytes,if_arg,if_present_arg,async_arg)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*), dimension(..), target, intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes
+    logical,intent(in),optional :: if_arg, if_present_arg
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    !
+    interface
+      subroutine gpufortrt_update_self_c_impl(hostptr,num_bytes,if_arg,if_present_arg) &
+              bind(c,name="gpufortrt_update_self")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: if_arg, if_present_arg
+      end subroutine
+      subroutine gpufortrt_update_self_async_c_impl(hostptr,num_bytes,if_arg,if_present_arg,async_arg) &
+              bind(c,name="gpufortrt_update_self_async")
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: if_arg, if_present_arg
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+    end interface
+    logical :: opt_if_arg, opt_if_present_arg
+    integer(c_size_t) :: opt_num_bytes
+    !
+    opt_if_arg = .true.
+    opt_if_present_arg = .false.
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(if_arg) ) opt_if_arg = if_arg
+    if ( present(if_present_arg) ) opt_if_present_arg = if_present_arg
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    !
+    if ( present(async_arg) ) then
+      call gpufortrt_update_self_async_c_impl(c_loc(hostptr),&
+                                                         opt_num_bytes,&
+                                                         logical(opt_if_arg,c_bool),&
+                                                         logical(opt_if_present_arg,c_bool),&
+                                                         async_arg)
+    else
+      call gpufortrt_update_self_c_impl(c_loc(hostptr),&
+                                                   opt_num_bytes,&
+                                                   logical(opt_if_arg,c_bool),&
+                                                   logical(opt_if_present_arg,c_bool))
+    endif
+  end subroutine
+
+  subroutine gpufortrt_update_device(hostptr,num_bytes,if_arg,if_present_arg,async_arg)
+    use iso_c_binding
+    use gpufortrt_types
+    implicit none
+    type(*),dimension(..),target,intent(in) :: hostptr
+    integer(c_size_t),intent(in),optional :: num_bytes
+    logical,intent(in),optional :: if_arg, if_present_arg
+    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
+    !
+    interface
+      subroutine gpufortrt_update_device_c_impl(hostptr,num_bytes,if_arg,if_present_arg) &
+              bind(c,name="gpufortrt_update_device")
+        use iso_c_binding
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: if_arg, if_present_arg
+      end subroutine
+      subroutine gpufortrt_update_device_async_c_impl(hostptr,num_bytes,if_arg,if_present_arg,async_arg) &
+              bind(c,name="gpufortrt_update_device_async")
+        use iso_c_binding
+        use gpufortrt_types
+        implicit none
+        type(c_ptr),value,intent(in) :: hostptr
+        integer(c_size_t),value,intent(in) :: num_bytes
+        logical(c_bool),value,intent(in) :: if_arg, if_present_arg
+        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
+      end subroutine
+    end interface
+    logical :: opt_if_arg, opt_if_present_arg
+    integer(c_size_t) :: opt_num_bytes
+    !
+    opt_if_arg = .true.
+    opt_if_present_arg = .false.
+    opt_num_bytes = int(sizeof(hostptr), kind=c_size_t)
+    if ( present(if_arg) ) opt_if_arg = if_arg
+    if ( present(if_present_arg) ) opt_if_present_arg = if_present_arg
+    if ( present(num_bytes) ) opt_num_bytes = num_bytes
+    !
+    if ( present(async_arg) ) then
+      call gpufortrt_update_device_async_c_impl(c_loc(hostptr),&
+                                                         opt_num_bytes,&
+                                                         logical(opt_if_arg,c_bool),&
+                                                         logical(opt_if_present_arg,c_bool),&
+                                                         async_arg)
+    else
+      call gpufortrt_update_device_c_impl(c_loc(hostptr),&
+                                                   opt_num_bytes,&
+                                                   logical(opt_if_arg,c_bool),&
+                                                   logical(opt_if_present_arg,c_bool))
+    endif
+  end subroutine
+end module
\ No newline at end of file
diff --git a/runtime/gpufortrt/src/gpufortrt_api.macros.f90 b/runtime/gpufortrt/src/gpufortrt_api.macros.f90
deleted file mode 100644
index 996ce01f..00000000
--- a/runtime/gpufortrt/src/gpufortrt_api.macros.f90
+++ /dev/null
@@ -1,496 +0,0 @@
-{# SPDX-License-Identifier: MIT #}
-{# Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved. #}
-{#######################################################################################}
-{% macro render_interface(routine,datatypes,max_rank,include_b_variant=True) %}
-{#######################################################################################}
-interface {{routine}}
-{% if include_b_variant %}
-  module procedure :: {{routine}}_b
-{% endif %}
-{% for triple in datatypes %}
-  module procedure :: {{routine}}0_{{triple[0]}}
-{%   for rank in range(1,max_rank+1) %}
-  module procedure :: {{routine}}{{rank}}_{{triple[0]}}
-{%   endfor %}{# rank #}
-{% endfor %}{# datatypes #}
-end interface
-{% endmacro %}
-{########################################################################################}
-{%- macro render_map_routines(datatypes,max_rank) -%}
-{# NOTE: type(*) is a Fortran 2018 feature.
-{########################################################################################}
-{% for clause in ["present","no_create","create","copyin","copy","copyout","delete"] -%}
-{%   set routine = "gpufortrt_map_" + clause %}
-function {{routine}}_b(hostptr,num_bytes,never_deallocate) result(retval)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  type(c_ptr),intent(in) :: hostptr
-  integer(c_size_t),intent(in),optional :: num_bytes 
-  logical,intent(in),optional :: never_deallocate
-  !
-  type(gpufortrt_mapping_t) :: retval
-  !
-  call gpufortrt_mapping_init(retval,hostptr,num_bytes,&
-         gpufortrt_map_kind_{{clause}},never_deallocate)
-end function
-
-{%   for triple in datatypes -%}
-!> \note never_deallocate only has effect on create,copyin,copyout, and copy mappings.
-function {{routine}}0_{{triple[0]}}(hostptr,never_deallocate) result(retval)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  logical,intent(in),optional :: never_deallocate
-  !
-  type(gpufortrt_mapping_t) :: retval
-  !
-  retval = {{routine}}_b(c_loc(hostptr),int({{triple[1]}},c_size_t),never_deallocate)
-end function
-
-{%     for rank in range(1,max_rank+1) %}
-!> \note never_deallocate only has effect on create,copyin,copyout, and copy mappings.
-function {{routine}}{{rank}}_{{triple[0]}}(hostptr,never_deallocate) result(retval)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  logical,intent(in),optional :: never_deallocate
-  !
-  type(gpufortrt_mapping_t) :: retval
-  !
-  retval = {{routine}}_b(c_loc(hostptr),int({{triple[1]}},c_size_t)*size(hostptr),never_deallocate)
-end function
-
-{%     endfor %} {# rank #}
-{%   endfor %} {# datatypes #}
-{% endfor %} {# clause #}
-{%- endmacro -%}
-{#######################################################################################}
-{% macro render_basic_delete_copyout_routines() %}
-{#######################################################################################}
-{% for clause in ["delete","copyout"] %}
-subroutine gpufortrt_{{clause}}_b(hostptr,num_bytes,async_arg,finalize)
-  use iso_c_binding
-  use gpufortrt_types, only: gpufortrt_handle_kind
-  implicit none
-  type(c_ptr), intent(in) :: hostptr
-  integer(c_size_t),intent(in),optional :: num_bytes
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  logical,intent(in),optional :: finalize
-  !
-  interface
-    subroutine gpufortrt_{{clause}}_c_impl(hostptr,num_bytes) &
-        bind(c,name="gpufortrt_{{clause}}")
-      use iso_c_binding
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-    end subroutine
-    subroutine gpufortrt_{{clause}}_finalize_c_impl(hostptr,num_bytes) &
-        bind(c,name="gpufortrt_{{clause}}_finalize")
-      use iso_c_binding
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-    end subroutine
-    subroutine gpufortrt_{{clause}}_async_c_impl(hostptr,num_bytes,async_arg) &
-        bind(c,name="gpufortrt_{{clause}}_async")
-      use iso_c_binding
-      use gpufortrt_types, only: gpufortrt_handle_kind
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-    end subroutine
-    subroutine gpufortrt_{{clause}}_finalize_async_c_impl(hostptr,num_bytes,async_arg) &
-        bind(c,name="gpufortrt_{{clause}}_finalize_async")
-      use iso_c_binding
-      use gpufortrt_types, only: gpufortrt_handle_kind
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-    end subroutine
-  end interface
-  !
-  if ( present(async_arg) ) then
-    if ( present(finalize) ) then
-      call gpufortrt_{{clause}}_finalize_async_c_impl(hostptr,num_bytes,async_arg)
-    else
-      call gpufortrt_{{clause}}_async_c_impl(hostptr,num_bytes,async_arg)
-    endif
-  else
-    if ( present(finalize) ) then
-      call gpufortrt_{{clause}}_finalize_c_impl(hostptr,num_bytes)
-    else
-      call gpufortrt_{{clause}}_c_impl(hostptr,num_bytes)
-    endif
-  endif
-end subroutine
-
-{% endfor %}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_specialized_delete_copyout_routines(datatypes,max_rank) %}
-{#######################################################################################}
-{% for clause in ["delete","copyout"] %}
-{%   for triple in datatypes -%}
-!> (Specialized version for Fortran scalar arguments)
-subroutine gpufortrt_{{clause}}0_{{triple[0]}}(hostptr,async_arg,finalize)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  logical,intent(in),optional :: finalize
-  !
-  call gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t),&
-                              async_arg,finalize)
-end subroutine
-
-{%     for rank in range(1,max_rank+1) %}
-!> (Specialized version for Fortran array arguments)
-subroutine gpufortrt_{{clause}}{{rank}}_{{triple[0]}}(hostptr,async_arg,finalize)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  logical,intent(in),optional :: finalize
-  !
-  call gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t)*size(hostptr),&
-                              async_arg,finalize)
-end subroutine
-
-{%     endfor %}{# rank #}
-{%   endfor %}{# datatypes #}
-{% endfor %}{# clause #}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_basic_copy_routines() %}
-{#######################################################################################}
-{% for clause in ["create","copyin","copy"] %}
-!> Map and directly return the corresponding deviceptr.
-function gpufortrt_{{clause}}_b(hostptr,num_bytes,never_deallocate,&
-    async_arg) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  type(c_ptr),intent(in) :: hostptr
-  integer(c_size_t),intent(in) :: num_bytes
-  logical,intent(in),optional :: never_deallocate
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg 
-  !
-  type(c_ptr) :: deviceptr
-  !
-  interface
-    function gpufortrt_{{clause}}_c_impl(hostptr,num_bytes,never_deallocate) &
-        bind(c,name="gpufortrt_{{clause}}") result(deviceptr)
-      use iso_c_binding
-      use gpufortrt_types
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      logical(c_bool),value,intent(in) :: never_deallocate
-      !
-      type(c_ptr) :: deviceptr
-    end function
-    function gpufortrt_{{clause}}_async_c_impl(hostptr,num_bytes,never_deallocate,async_arg) &
-        bind(c,name="gpufortrt_{{clause}}_async") result(deviceptr)
-      use iso_c_binding
-      use gpufortrt_types
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      logical(c_bool),value,intent(in) :: never_deallocate
-      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-      !
-      type(c_ptr) :: deviceptr
-    end function
-  end interface
-  !
-  logical(c_bool) :: opt_never_deallocate
-  !
-  opt_never_deallocate = .false._c_bool
-  if ( present(never_deallocate) ) opt_never_deallocate = never_deallocate
-  if ( present(async_arg) ) then
-    deviceptr = gpufortrt_{{clause}}_async_c_impl(hostptr,num_bytes,opt_never_deallocate,async_arg)
-  else
-    deviceptr = gpufortrt_{{clause}}_c_impl(hostptr,num_bytes,opt_never_deallocate)
-  endif
-end function
-
-{% endfor %}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_specialized_copy_routines(datatypes,max_rank) %}
-{#######################################################################################}
-{% for clause in ["create","copyin","copy"] %}
-{%   for triple in datatypes -%}
-!> Map and directly return the corresponding deviceptr.
-!> (Specialized version for Fortran scalar arguments)
-function gpufortrt_{{clause}}0_{{triple[0]}}(hostptr,never_deallocate,async_arg) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  logical,intent(in),optional :: never_deallocate
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg 
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t),&
-                                     never_deallocate,async_arg)
-end function
-
-{%     for rank in range(1,max_rank+1) %}
-!> Map and directly return the corresponding deviceptr.
-!> (Specialized version for Fortran array arguments)
-function gpufortrt_{{clause}}{{rank}}_{{triple[0]}}(hostptr,&
-    never_deallocate,async_arg) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  logical,intent(in),optional :: never_deallocate
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg 
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t)*size(hostptr),&
-                                     never_deallocate,async_arg)
-end function
-{%     endfor %}{# rank #}
-{%   endfor %}{# datatypes #}
-{% endfor %}{# clause #}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_specialized_deviceptr_routines(datatypes,max_rank) %}
-{#######################################################################################}
-{% for triple in datatypes -%}
-function gpufortrt_deviceptr0_{{triple[0]}}(hostptr) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_deviceptr_b(c_loc(hostptr))
-end function
-
-{%   for rank in range(1,max_rank+1) %}
-function gpufortrt_deviceptr{{rank}}_{{triple[0]}}(hostptr) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_deviceptr_b(c_loc(hostptr))
-end function
-
-{%   endfor %}{# rank #}
-{% endfor %}{# datatypes #}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_specialized_present_routines(datatypes,max_rank) %}
-{#######################################################################################}
-{% for clause in ["present"] %}
-{%   for triple in datatypes -%}
-function gpufortrt_{{clause}}0_{{triple[0]}}(hostptr) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t))
-end function
-
-{%     for rank in range(1,max_rank+1) %}
-function gpufortrt_{{clause}}{{rank}}_{{triple[0]}}(hostptr) result(deviceptr)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  !
-  type(c_ptr) :: deviceptr
-  !
-  deviceptr = gpufortrt_{{clause}}_b(c_loc(hostptr),int({{triple[1]}},kind=c_size_t)*size(hostptr))
-end function
-
-{%     endfor %}{# rank #}
-{%   endfor %}{# datatypes #}
-{% endfor %}{# clause #}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_map_interfaces(datatypes,max_rank) %}
-{#######################################################################################}
-{% for clause in ["present","no_create","create","copyin","copy","copyout","delete"] %}
-{{ render_interface("gpufortrt_map_"+clause,datatypes,max_rank) }}
-{% endfor %}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_copy_interfaces(datatypes,max_rank) %}
-{#######################################################################################}
-{% for clause in ["present"] %}
-{{ render_interface("gpufortrt_"+clause,datatypes,max_rank,False) }}
-
-{% endfor %}
-{% for clause in ["create","copyin","copy"] %}
-{{ render_interface("gpufortrt_"+clause,datatypes,max_rank) }}
-
-{% endfor %}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_basic_update_routine(update_kind) %}
-{#######################################################################################}
-subroutine gpufortrt_update_{{update_kind}}_b(hostptr,num_bytes,if_arg,if_present_arg,async_arg)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  type(c_ptr),intent(in) :: hostptr
-  integer(c_size_t),intent(in) :: num_bytes
-  logical,intent(in),optional :: if_arg, if_present_arg
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  !
-  interface
-    subroutine gpufortrt_update_{{update_kind}}_c_impl(hostptr,num_bytes,if_arg,if_present_arg) &
-            bind(c,name="gpufortrt_update_{{update_kind}}")
-      use iso_c_binding
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      logical(c_bool),value,intent(in) :: if_arg, if_present_arg
-    end subroutine
-    subroutine gpufortrt_update_{{update_kind}}_async_c_impl(hostptr,num_bytes,if_arg,if_present_arg,async_arg) &
-            bind(c,name="gpufortrt_update_{{update_kind}}_async")
-      use iso_c_binding
-      use gpufortrt_types
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      logical(c_bool),value,intent(in) :: if_arg, if_present_arg
-      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-    end subroutine
-  end interface
-  logical :: opt_if_arg, opt_if_present_arg
-  !
-  opt_if_arg = .true.
-  opt_if_present_arg = .false.
-  if ( present(if_arg) ) opt_if_arg = if_arg
-  if ( present(if_present_arg) ) opt_if_present_arg = if_present_arg
-  !
-  if ( present(async_arg) ) then
-    call gpufortrt_update_{{update_kind}}_async_c_impl(hostptr,&
-                                                       num_bytes,&
-                                                       logical(opt_if_arg,c_bool),&
-                                                       logical(opt_if_present_arg,c_bool),&
-                                                       async_arg)
-  else
-    call gpufortrt_update_{{update_kind}}_c_impl(hostptr,&
-                                                 num_bytes,&
-                                                 logical(opt_if_arg,c_bool),&
-                                                 logical(opt_if_present_arg,c_bool))
-  endif
-end subroutine
-{% endmacro %}
-{#######################################################################################}
-{% macro render_specialized_update_routines(update_kind,datatypes,max_rank) %}
-{#######################################################################################}
-{% for triple in datatypes %}
-subroutine gpufortrt_update_{{update_kind}}0_{{triple[0]}}(hostptr,if_arg,if_present_arg,async_arg)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  logical,intent(in),optional :: if_arg, if_present_arg
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  !
-  call gpufortrt_update_{{update_kind}}_b(c_loc(hostptr),int({{triple[1]}},c_size_t),if_arg,if_present_arg,async_arg)
-end subroutine
-
-{%   for rank in range(1,max_rank+1) %}
-subroutine gpufortrt_update_{{update_kind}}{{rank}}_{{triple[0]}}(hostptr,if_arg,if_present_arg,async_arg)
-  use iso_c_binding
-  use gpufortrt_types
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr(:{% for i in range(1,rank) %},:{% endfor %})
-  logical,intent(in),optional :: if_arg, if_present_arg
-  integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-  !
-  call gpufortrt_update_{{update_kind}}_b(c_loc(hostptr),int({{triple[1]}},c_size_t)*size(hostptr),if_arg,if_present_arg,async_arg)
-end subroutine
-
-{%   endfor %}{# rank #}
-{% endfor %}{# datatypes #}
-{% endmacro %}
-{########################################################################################}
-{%- macro render_set_fptr_lower_bound(fptr,rank) -%}
-{########################################################################################}
-{{fptr}}(&
-{% for i in range(1,rank+1) %}
-  opt_lbounds({{i}}):{{ "," if not loop.last else ")" }}&
-{% endfor %}
-    => {{fptr}}
-{%- endmacro  -%}
-{#######################################################################################}
-{% macro render_specialized_use_device_routines(datatypes,max_rank) %}
-{#######################################################################################}
-{% for triple in datatypes %}
-subroutine gpufortrt_use_device0_{{triple[0]}}(resultptr,hostptr,if_arg,if_present_arg)
-  use iso_c_binding
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr
-  logical,intent(in),optional :: if_arg, if_present_arg
-  !
-  {{triple[2]}},pointer,intent(inout) :: resultptr
-  !
-  type(c_ptr) :: tmp_cptr
-  !
-  tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
-  call c_f_pointer(tmp_cptr,resultptr)
-end subroutine
-
-{% for rank in range(1,max_rank+1) %}
-subroutine gpufortrt_use_device{{rank}}_{{triple[0]}}(resultptr,hostptr,sizes,lbounds,if_arg,if_present_arg)
-  use iso_c_binding
-  implicit none
-  {{triple[2]}},target,intent(in) :: hostptr({% for i in range(1,rank) %}1:1,{% endfor %}*)
-  integer,intent(in),optional :: sizes({{rank}}), lbounds({{rank}})
-  logical,intent(in),optional :: if_arg, if_present_arg
-  !
-  {{triple[2]}},pointer,intent(inout) :: resultptr(:{% for i in range(1,rank) %},:{% endfor %})
-  !
-  type(c_ptr) :: tmp_cptr
-  integer :: opt_sizes({{rank}}), opt_lbounds({{rank}})
-  !
-  opt_sizes = 1
-  opt_lbounds = 1  
-  if ( present(sizes) ) opt_sizes = sizes
-  if ( present(lbounds) ) opt_lbounds = lbounds
-  tmp_cptr = gpufortrt_use_device_b(c_loc(hostptr),if_arg,if_present_arg)
-  call c_f_pointer(tmp_cptr,resultptr,shape=opt_sizes)
-{{ render_set_fptr_lower_bound("resultptr",rank) | indent(2,True) }}
-end subroutine
-
-{%   endfor %}{# rank #}
-{% endfor %}{# datatypes #}
-{% endmacro %}
-{#######################################################################################}
-{% macro render_use_device_interface(datatypes,max_rank) %}
-{#######################################################################################}
-interface gpufortrt_use_device
-{% for rank in range(0,max_rank+1) %}
-{%   for triple in datatypes %}
-  module procedure :: gpufortrt_use_device{{rank}}_{{triple[0]}}
-{%   endfor %}{# datatypes #}
-{% endfor %}{# rank #}
-end interface
-{% endmacro %}
-{#######################################################################################}
diff --git a/runtime/gpufortrt/src/gpufortrt_api.template.f90 b/runtime/gpufortrt/src/gpufortrt_api.template.f90
deleted file mode 100644
index cc58e139..00000000
--- a/runtime/gpufortrt/src/gpufortrt_api.template.f90
+++ /dev/null
@@ -1,42 +0,0 @@
-{# SPDX-License-Identifier: MIT #}
-{# Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved. #}
-{########################################################################################}
-{% import "src/gpufortrt_api.macros.f90" as gam %}
-{########################################################################################}
-! SPDX-License-Identifier: MIT
-! Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
-module gpufortrt_api
-  use gpufortrt_api_core
-{{ gam.render_map_interfaces(datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_copy_interfaces(datatypes,max_rank) | indent(2,True) }}
-{{ gam.render_interface("gpufortrt_delete",datatypes,max_rank) | indent(2,True) }}
-{{ gam.render_interface("gpufortrt_copyout",datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_interface("gpufortrt_update_self",datatypes,max_rank) | indent(2,True) }}
-{{ gam.render_interface("gpufortrt_update_device",datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_use_device_interface(datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_interface("gpufortrt_deviceptr",datatypes,max_rank,False) | indent(2,True) }}
-
-contains
-
-{{ gam.render_map_routines(datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_basic_copy_routines() | indent(2,True) }}
-{{ gam.render_specialized_copy_routines(datatypes,max_rank) | indent(2,True) }}
-{{ gam.render_specialized_present_routines(datatypes,max_rank) | indent(2,True) }}
-{{ gam.render_basic_delete_copyout_routines() | indent(2,True) }}
-{{ gam.render_specialized_delete_copyout_routines(datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_basic_update_routine("self") }}
-{{ gam.render_specialized_update_routines("self",datatypes,max_rank) }}
-
-{{ gam.render_basic_update_routine("device") | indent(2,True) }}
-{{ gam.render_specialized_update_routines("device",datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_specialized_use_device_routines(datatypes,max_rank) | indent(2,True) }}
-
-{{ gam.render_specialized_deviceptr_routines(datatypes,max_rank) | indent(2,True) }}
-end module
diff --git a/runtime/gpufortrt/src/gpufortrt_api_core.f90 b/runtime/gpufortrt/src/gpufortrt_api_core.f90
deleted file mode 100644
index 9af53d50..00000000
--- a/runtime/gpufortrt/src/gpufortrt_api_core.f90
+++ /dev/null
@@ -1,295 +0,0 @@
-! SPDX-License-Identifier: MIT
-! Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
-module gpufortrt_api_core
-  use gpufortrt_types
-
-  interface
-    subroutine gpufortrt_init() bind(c,name="gpufortrt_init")
-      implicit none
-    end subroutine
-  
-    subroutine gpufortrt_shutdown() bind(c,name="gpufortrt_shutdown")
-      implicit none
-    end subroutine
-  
-    function gpufortrt_get_stream(async_arg) &
-        bind(c,name="gpufortrt_get_stream") &
-          result(stream)
-      use iso_c_binding, only: c_ptr
-      use gpufortrt_types, only: gpufortrt_handle_kind
-      implicit none
-      integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-      !
-      type(c_ptr) :: stream
-    end function
-  end interface
-   
-  interface gpufortrt_present
-    function gpufortrt_present_b(hostptr,num_bytes) &
-        bind(c,name="gpufortrt_present") result(deviceptr)
-      use iso_c_binding, only: c_ptr, c_size_t
-      use gpufortrt_types, only: gpufortrt_counter_none
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      integer(c_size_t),value,intent(in) :: num_bytes
-      !
-      type(c_ptr) :: deviceptr
-    end function
-  end interface
-
-  interface gpufortrt_deviceptr
-    function gpufortrt_deviceptr_b(hostptr) & 
-        bind(c,name="gpufortrt_deviceptr") &
-          result(deviceptr)
-      use iso_c_binding, only: c_ptr, c_size_t
-      implicit none
-      type(c_ptr),value,intent(in) :: hostptr
-      !
-      type(c_ptr) :: deviceptr
-    end function
-  end interface
-
-contains
-  
-  !> Ignore the result of a mapping routine.
-  !> \param[in] deviceptr a device pointer.
-  subroutine gpufortrt_ignore(deviceptr)
-    use iso_c_binding, only: c_ptr
-    implicit none
-    type(c_ptr),intent(in) :: deviceptr
-    ! nop
-  end subroutine
-
-  subroutine gpufortrt_wait(wait_arg,async_arg,condition)
-    use iso_c_binding
-    implicit none
-    integer(gpufortrt_handle_kind),dimension(:),target,intent(in),optional :: wait_arg,async_arg
-    logical,intent(in),optional :: condition
-    !
-    interface
-      subroutine gpufortrt_wait_all_c_impl(condition) &
-          bind(c,name="gpufortrt_wait_all")
-        use iso_c_binding
-        implicit none
-        logical(c_bool),value,intent(in):: condition 
-      end subroutine
-      subroutine gpufortrt_wait_all_async_c_impl(async_arg,num_async_args,condition) &
-          bind(c,name="gpufortrt_wait_all_async")
-        use iso_c_binding
-        implicit none
-        type(c_ptr),value,intent(in) :: async_arg
-        integer(c_int),value,intent(in) :: num_async_args
-        logical(c_bool),value,intent(in) :: condition 
-      end subroutine
-      subroutine gpufortrt_wait_c_impl(wait_arg,num_wait_args,condition) &
-          bind(c,name="gpufortrt_wait")
-        use iso_c_binding
-        implicit none
-        type(c_ptr),value,intent(in) :: wait_arg
-        integer(c_int),value,intent(in) :: num_wait_args
-        logical(c_bool),value,intent(in) :: condition 
-      end subroutine
-      subroutine gpufortrt_wait_async_c_impl(wait_arg,num_wait_args,&
-                                             async_arg,num_async_args,&
-                                             condition) &
-          bind(c,name="gpufortrt_wait_async")
-        use iso_c_binding
-        implicit none
-        type(c_ptr),value,intent(in) :: wait_arg, async_arg
-        integer(c_int),value,intent(in) :: num_wait_args, num_async_args
-        logical(c_bool),value,intent(in) :: condition 
-      end subroutine
-    end interface
-    !
-    logical(c_bool) :: opt_if_arg
-    !
-    opt_if_arg = .true._c_bool
-    if ( present(condition) ) opt_if_arg = logical(condition,kind=c_bool)
-    !
-    if ( present(wait_arg) ) then
-      if ( present(async_arg) ) then
-        call gpufortrt_wait_async_c_impl(&
-             c_loc(wait_arg),size(wait_arg,kind=c_int),&
-             c_loc(async_arg),size(async_arg,kind=c_int),&
-             opt_if_arg)
-      else
-        call gpufortrt_wait_c_impl(&
-             c_loc(wait_arg),size(wait_arg,kind=c_int),&
-             opt_if_arg)
-      endif
-    else
-      if ( present(async_arg) ) then
-        call gpufortrt_wait_all_async_c_impl(&
-             c_loc(async_arg),size(async_arg,kind=c_int),&
-             opt_if_arg)
-      else
-        call gpufortrt_wait_all_c_impl(opt_if_arg)
-      endif
-    endif
-  end subroutine
-
-  subroutine gpufortrt_data_start(mappings,async_arg)
-  !subroutine gpufortrt_data_start(device_kind,mappings,async_arg)
-    use iso_c_binding
-    implicit none
-    !integer,intent(in) :: device_kind
-    type(gpufortrt_mapping_t),dimension(:),target,intent(in),optional :: mappings
-    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-    !
-    interface
-      subroutine gpufortrt_data_start_c_impl(mappings,num_mappings) bind(c,name="gpufortrt_data_start")
-        use iso_c_binding
-        implicit none
-        type(c_ptr),intent(in),value :: mappings
-        integer(c_int),intent(in),value :: num_mappings
-      end subroutine
-      subroutine gpufortrt_data_start_async_c_impl(mappings,num_mappings,async_arg) bind(c,name="gpufortrt_data_start_async")
-        use iso_c_binding
-        use gpufortrt_types, only: gpufortrt_handle_kind
-        implicit none
-        type(c_ptr),intent(in),value :: mappings
-        integer(c_int),intent(in),value :: num_mappings
-        integer(gpufortrt_handle_kind),intent(in),value :: async_arg
-      end subroutine
-    end interface 
-    !
-    if ( present(async_arg) ) then
-      if ( present(mappings) ) then
-        call gpufortrt_data_start_async_c_impl(c_loc(mappings),size(mappings),&
-                                               int(async_arg,kind=c_int))
-      else
-        call gpufortrt_data_start_async_c_impl(c_null_ptr,0_c_int,&
-                                               int(async_arg,kind=c_int)) 
-      endif
-    else
-      if ( present(mappings) ) then
-        call gpufortrt_data_start_c_impl(c_loc(mappings),size(mappings))  
-      else
-        call gpufortrt_data_start_c_impl(c_null_ptr,0_c_int)  
-      endif
-    endif
-  end subroutine
-   
-  subroutine gpufortrt_data_end(async_arg)
-    implicit none
-    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-    !
-    interface
-      subroutine gpufortrt_data_end_c_impl() bind(c,name="gpufortrt_data_end")
-        use iso_c_binding
-        use gpufortrt_types, only: gpufortrt_handle_kind
-        implicit none
-      end subroutine
-      subroutine gpufortrt_data_end_async_c_impl(async_arg) bind(c,name="gpufortrt_data_end_async")
-        use iso_c_binding
-        use gpufortrt_types, only: gpufortrt_handle_kind
-        implicit none
-        integer(gpufortrt_handle_kind),intent(in),value :: async_arg
-      end subroutine
-    end interface 
-    !
-    if ( present(async_arg) ) then
-      call gpufortrt_data_end_async_c_impl(async_arg)  
-    else
-      call gpufortrt_data_end_c_impl()  
-    endif
-  end subroutine
-
-  subroutine gpufortrt_enter_exit_data(mappings,async_arg,finalize)
-    use iso_c_binding
-    implicit none
-    !
-    !integer,intent(in) :: device_kind
-    type(gpufortrt_mapping_t),dimension(:),target,intent(in),optional :: mappings
-    integer(gpufortrt_handle_kind),intent(in),optional :: async_arg
-    logical,intent(in),optional :: finalize
-    !
-    interface
-      subroutine gpufortrt_enter_exit_data_c_impl(mappings,num_mappings,finalize) bind(c,name="gpufortrt_enter_exit_data")
-        use iso_c_binding
-        implicit none
-        type(c_ptr),value,intent(in) :: mappings
-        integer(c_int),value,intent(in) :: num_mappings
-        logical(c_bool),value,intent(in) :: finalize
-      end subroutine
-      subroutine gpufortrt_enter_exit_data_async_c_impl(mappings,num_mappings,async_arg,finalize) bind(c,name="gpufortrt_enter_exit_data_async")
-        use iso_c_binding
-        use gpufortrt_types, only: gpufortrt_handle_kind
-        implicit none
-        type(c_ptr),value,intent(in) :: mappings
-        integer(c_int),value,intent(in) :: num_mappings
-        integer(gpufortrt_handle_kind),value,intent(in) :: async_arg
-        logical(c_bool),value,intent(in) :: finalize
-      end subroutine
-    end interface 
-    !
-    logical(c_bool) :: opt_finalize
-    !
-    opt_finalize = .false._c_bool
-    if ( present(finalize) ) opt_finalize = logical(finalize,kind=c_bool)
-    !
-    if ( present(async_arg) ) then
-      if ( present(mappings) ) then
-        call gpufortrt_enter_exit_data_async_c_impl(&
-          c_loc(mappings),&
-          size(mappings,kind=c_int),&
-          async_arg,&
-          opt_finalize) 
-      else
-        call gpufortrt_enter_exit_data_async_c_impl(&
-          c_null_ptr,&
-          0_c_int,&
-          async_arg,&
-          opt_finalize) 
-      endif
-    else
-      if ( present(mappings) ) then
-        call gpufortrt_enter_exit_data_c_impl(&
-          c_loc(mappings),&
-          size(mappings,kind=c_int),&
-          opt_finalize) 
-      else
-        call gpufortrt_enter_exit_data_c_impl(&
-          c_null_ptr,&
-          0_c_int,&
-          opt_finalize) 
-      endif
-    endif
-  end subroutine
-
-  !> Lookup device pointer for given host pointer.
-  !> \param[in] condition condition that must be met, otherwise host pointer is returned. Defaults to '.true.'.
-  !> \param[in] if_present Only return device pointer if one could be found for the host pointer.
-  !>                       otherwise host pointer is returned. Defaults to '.false.'.
-  !> \note Returns a c_null_ptr if the host pointer is invalid, i.e. not C associated.
-  function gpufortrt_use_device_b(hostptr,condition,if_present) result(resultptr)
-    use iso_c_binding
-    implicit none
-    type(c_ptr),intent(in) :: hostptr
-    logical,intent(in),optional :: condition, if_present
-    !
-    type(c_ptr) :: resultptr
-    !
-    interface
-      function gpufortrt_use_device_c_impl(hostptr,condition,if_present) &
-          bind(c,name="gpufortrt_use_device") result(deviceptr)
-        use iso_c_binding
-        implicit none
-        type(c_ptr),value,intent(in) :: hostptr
-        logical(c_bool),value,intent(in) :: condition, if_present
-        !
-        type(c_ptr) :: deviceptr
-      end function
-    end interface
-    !
-    logical(c_bool) :: opt_if_arg, opt_if_present_arg
-    !
-    opt_if_arg = .true._c_bool
-    opt_if_present_arg = .false._c_bool
-    if ( present(condition) ) opt_if_arg = logical(condition,kind=c_bool)
-    if ( present(if_present) ) opt_if_present_arg = logical(if_present,kind=c_bool)
-    !
-    resultptr = gpufortrt_use_device_c_impl(hostptr,opt_if_arg,opt_if_present_arg)
-  end function
-
-end module
diff --git a/runtime/gpufortrt/src/internal/gpufortrt_core.h b/runtime/gpufortrt/src/internal/gpufortrt_core.h
index 53eaec78..80e11418 100644
--- a/runtime/gpufortrt/src/internal/gpufortrt_core.h
+++ b/runtime/gpufortrt/src/internal/gpufortrt_core.h
@@ -353,6 +353,7 @@ namespace gpufortrt {
       queue_record_t& operator[](const int i);
       const queue_record_t& operator[](const int i) const;
      
+      size_t size();
       /** Reserve space for `capacity` queues. */
       void reserve(const int capacity);
 
diff --git a/runtime/gpufortrt/src/internal/queue_record_list_t.cpp b/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
index 8bce782a..639fd5b6 100644
--- a/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
+++ b/runtime/gpufortrt/src/internal/queue_record_list_t.cpp
@@ -18,6 +18,10 @@ const gpufortrt::internal::queue_record_t& gpufortrt::internal::queue_record_lis
   return this->records[i];  
 }
 
+size_t gpufortrt::internal::queue_record_list_t::size() {
+  return this->records.size();
+}
+
 void gpufortrt::internal::queue_record_list_t::reserve(int capacity) {
   this->records.reserve(capacity); 
 }
@@ -70,13 +74,13 @@ gpufortrt_queue_t gpufortrt::internal::queue_record_list_t::use_create_queue(int
 namespace {
   void synchronize_default_queue() {
     gpufortrt::internal::queue_record_t default_queue_record;
-    default_queue_record.setup(-1,gpufortrt_default_queue);
+    default_queue_record.setup(-1); // original was (-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);
+    default_queue_record.setup(-1); // original was (-1,gpufortrt_default_queue)
     return default_queue_record.test();
   }
 }
@@ -110,7 +114,7 @@ void gpufortrt::internal::queue_record_list_t::synchronize(const int id) {
     std::size_t loc;
     std::tie(success,loc) = this->find_record(id);
     if ( success ) {
-      result = this->records[loc].synchronize(); 
+      this->records[loc].synchronize(); 
       LOG_INFO(3,"synchronize queue; " << this->records[loc])
     } else {
       LOG_INFO(3,"synchronize queue; no queue found for id="<<id)
@@ -125,7 +129,7 @@ void gpufortrt::internal::queue_record_list_t::synchronize(const int id) {
  * 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++) {
+  for (size_t i = 0; i < gpufortrt::internal::queue_record_list.records.size(); i++) {
     auto& queue_record = gpufortrt::internal::queue_record_list[i];
     result &= queue_record.test();
   } 
@@ -137,7 +141,7 @@ bool gpufortrt::internal::queue_record_list_t::test_all() {
 /** 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++) {
+  for (size_t i = 0; i < gpufortrt::internal::queue_record_list.records.size(); i++) {
     auto& queue_record = gpufortrt::internal::queue_record_list[i];
     queue_record.synchronize();
   } 
diff --git a/runtime/gpufortrt/src/internal/queue_record_t.cpp b/runtime/gpufortrt/src/internal/queue_record_t.cpp
index 0569e323..940c6a0d 100644
--- a/runtime/gpufortrt/src/internal/queue_record_t.cpp
+++ b/runtime/gpufortrt/src/internal/queue_record_t.cpp
@@ -43,7 +43,7 @@ void gpufortrt::internal::queue_record_t::destroy() {
 
 bool gpufortrt::internal::queue_record_t::test() {
   bool result = hipStreamQuery(this->queue) == hipSuccess;
-  LOG_INFO(4,"<test queue; " << *this" <<; result:" << result << )
+  LOG_INFO(4,"<test queue; " << *this <<" ; result:" << result )
   return result;
 }
 
diff --git a/runtime/gpufortrt/src/obj.mk b/runtime/gpufortrt/src/obj.mk
index 08ed9dee..15b4399d 100644
--- a/runtime/gpufortrt/src/obj.mk
+++ b/runtime/gpufortrt/src/obj.mk
@@ -1,8 +1,10 @@
 # SPDX-License-Identifier: MIT
 # Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
 F_OBJ += ./gpufortrt_types.o \
-         ./gpufortrt_api_core.o \
-         ./gpufortrt_api.o
+         ./gpufortrt_api.o \
+         ./openacc.o
+        #  ./gpufortrt_api_core.o \
 
 CXX_OBJ += ./gpufortrt_types.cpp.o \
-		   ./gpufortrt_api.cpp.o
+		   ./gpufortrt_api.cpp.o \
+           ./openacc.cpp.o
diff --git a/runtime/gpufortrt/src/openacc.cpp b/runtime/gpufortrt/src/openacc.cpp
index 54114188..719b7875 100644
--- a/runtime/gpufortrt/src/openacc.cpp
+++ b/runtime/gpufortrt/src/openacc.cpp
@@ -11,17 +11,34 @@ namespace {
 
   acc_device_t current_device = acc_device_none;
 
-  bool check_device_type(acc_device_t dev_type,bool allow_none=true) {
+  gpufortrt_device_property_t check_device_property(acc_device_property_t property,bool allow_none=true){
+    switch(property){
+      case acc_property_memory:
+        return gpufortrt_property_memory;
+      case acc_property_free_memory:
+        return gpufortrt_property_free_memory;
+      case acc_property_shared_memory_support:
+        return gpufortrt_property_shared_memory_support;
+      case acc_property_name:
+        return gpufortrt_property_name;
+      case acc_property_vendor:
+        return gpufortrt_property_vendor;
+      case acc_property_driver:
+        return gpufortrt_property_driver;
+    }
+  }
+
+  acc_device_t 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_default:
+        return ::default_device;
       case acc_device_current:
-        return ::check_device_type(::current_device);
+        return ::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_not_host: 
+      case acc_device_hip: 
+      case acc_device_radeon: 
       case acc_device_nvidia:
         return acc_device_hip;
       case acc_device_none:
@@ -46,7 +63,7 @@ acc_device_t acc_get_device_type(void) {
 
 int acc_get_num_devices(acc_device_t dev_type) {
   if ( ::check_device_type(dev_type) == acc_device_hip ) {
-    return gpufortrt_get_device_num();
+    return gpufortrt_get_num_devices();
   } else { // host
     return 1;
   }
@@ -77,7 +94,7 @@ 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);
+    return gpufortrt_get_property(dev_num,check_device_property(property));
   } else {
     throw std::invalid_argument("acc_get_property: only implemented for non-host device types");
   }
@@ -87,13 +104,34 @@ 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);
+    return gpufortrt_get_property_string(dev_num,check_device_property(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
+size_t acc_get_property_f(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_f(dev_num,check_device_property(property));
+  } else {
+    throw std::invalid_argument("acc_get_property: only implemented for non-host device types");
+  }
+}
+
+const 
+char* acc_get_property_string_f(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_f(dev_num,check_device_property(property));
+  } else {
+    throw std::invalid_argument("acc_get_property_string: only implemented for non-host device types");
+  }
+}
+
 void acc_set_device_num_f(int dev_num, acc_device_t dev_type) {
   acc_set_device_num(dev_num-1,dev_type);
 }
@@ -102,7 +140,7 @@ int acc_get_device_num_f(acc_device_t dev_type) {
 }
 
 
-void acc_init(acc_on_device_t dev_type) {
+void acc_init(acc_device_t dev_type) {
   if ( ::check_device_type(dev_type) == acc_device_hip ) {
     gpufortrt_init();
   }
@@ -113,3 +151,152 @@ void acc_shutdown(acc_device_t dev_type) {
     gpufortrt_shutdown();
   }
 }
+
+int acc_is_present(void* hostptr,std::size_t num_bytes) {
+  if(gpufortrt_is_present(hostptr, num_bytes)) 
+    return 1;
+  else 
+    return 0;
+}
+
+void acc_wait(int wait_arg){
+  gpufortrt_wait(&wait_arg, 1, true);
+}
+
+void acc_wait_async(int wait_arg, int async_arg){
+  gpufortrt_wait_async(&wait_arg, 1, &async_arg, 1, true);
+}
+
+void acc_wait_all(){
+  gpufortrt_wait_all(true);
+}
+
+void acc_wait_all_async(int async_arg){
+  gpufortrt_wait_all_async(&async_arg, 1, true);
+}
+
+void acc_wait_device(int wait_arg, int dev_num){
+  gpufortrt_wait_device(&wait_arg, 1, dev_num, true);
+}
+
+void acc_wait_device_async(int wait_arg, int async_arg, int dev_num){
+  gpufortrt_wait_device_async(&wait_arg, 1, &async_arg, 1, dev_num, true);
+}
+
+void acc_wait_all_device(int dev_num){
+  gpufortrt_wait_all_device(dev_num, true);
+}
+
+void acc_wait_all_device_async(int async_arg, int dev_num){
+  gpufortrt_wait_all_device_async(&async_arg, 1, dev_num, true);
+}
+
+int acc_async_test(int wait_arg){
+  return gpufortrt_async_test(wait_arg);
+}
+
+int acc_async_test_device(int wait_arg, int dev_num){
+  return gpufortrt_async_test_device( wait_arg, dev_num);
+}
+
+int acc_async_test_all(void){
+  return gpufortrt_async_test_all();
+}
+
+int acc_async_test_all_device(int dev_num){
+  return gpufortrt_async_test_all_device(dev_num);
+}
+
+d_void* acc_copyin(h_void* data_arg, size_t bytes){
+  return gpufortrt_copyin(data_arg, bytes, true);
+}
+
+void acc_copyin_async(h_void* data_arg, size_t bytes,
+                      int async_arg){
+  gpufortrt_copyin_async(data_arg, bytes, async_arg, false);
+}
+
+d_void* acc_create(h_void* data_arg, size_t bytes){
+  return gpufortrt_create(data_arg, bytes, false);
+}
+
+void acc_create_async(h_void* data_arg, size_t bytes,
+                      int async_arg){
+  gpufortrt_create_async(data_arg, bytes, async_arg, false);
+}
+
+void acc_copyout(h_void* data_arg, size_t bytes){
+  gpufortrt_copyout(data_arg, bytes);
+}
+void acc_copyout_async(h_void* data_arg, size_t bytes, 
+                        int async_arg){
+  gpufortrt_copyout_async(data_arg, bytes, async_arg);
+}
+void acc_copyout_finalize(h_void* data_arg, size_t bytes){
+  gpufortrt_copyout_finalize(data_arg, bytes);
+}
+void acc_copyout_finalize_async(h_void* data_arg, size_t bytes, 
+                                int async_arg){
+  gpufortrt_copyout_finalize_async( data_arg, bytes, async_arg);
+}
+
+int acc_get_default_async(void){
+  return gpufortrt_get_default_async();
+}
+
+void acc_set_default_async(int async_arg){
+  gpufortrt_set_default_async(async_arg);
+}
+
+void acc_delete(h_void* data_arg, size_t bytes){
+  gpufortrt_delete( data_arg, bytes);
+}
+
+void acc_delete_async(h_void* data_arg, size_t bytes,
+                      int async_arg){
+  gpufortrt_delete_async(data_arg, bytes, async_arg);
+}
+
+void acc_delete_finalize(h_void* data_arg, size_t bytes){
+  gpufortrt_delete_finalize( data_arg, bytes);
+}
+
+void acc_delete_finalize_async(h_void* data_arg,
+                               size_t bytes, int async_arg){
+  gpufortrt_delete_finalize_async(data_arg, bytes, async_arg);
+}
+
+void acc_update_device(h_void* data_arg, size_t bytes){
+  gpufortrt_update_device(data_arg, bytes, true, false);
+}
+
+void acc_update_device_async(h_void* data_arg, size_t bytes,
+                             int async_arg){
+  gpufortrt_update_device_async(data_arg, bytes, true, false, async_arg);
+}
+
+void acc_update_self(h_void* data_arg, size_t bytes){
+  gpufortrt_update_self(data_arg, bytes, true, false);
+}
+
+void acc_update_self_async(h_void* data_arg, size_t bytes,
+                           int async_arg){
+  gpufortrt_update_self_async(data_arg, bytes, true, false, async_arg);
+}
+
+d_void* acc_deviceptr(h_void* data_arg){
+  return gpufortrt_deviceptr(data_arg);
+}
+
+d_void* acc_malloc(size_t bytes){ 
+  return gpufortrt_malloc(bytes);
+}
+
+void acc_free(d_void* data_dev){ 
+  gpufortrt_free(data_dev);
+}
+
+void acc_map_data(h_void* data_arg, d_void* data_dev,
+                  size_t bytes){
+  gpufortrt_map_data(data_arg, data_dev, bytes);
+}
\ No newline at end of file
diff --git a/runtime/gpufortrt/src/openacc.f90 b/runtime/gpufortrt/src/openacc.f90
index 59c507a6..e6396569 100644
--- a/runtime/gpufortrt/src/openacc.f90
+++ b/runtime/gpufortrt/src/openacc.f90
@@ -10,16 +10,15 @@ module openacc
   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
+    enumerator :: acc_device_hip 
+    enumerator :: acc_device_radeon 
+    enumerator :: acc_device_nvidia 
   end enum
   
   integer, parameter :: acc_device_kind = kind(acc_device_none)
@@ -36,626 +35,983 @@ module openacc
   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 
 
+  interface acc_is_present
+    module procedure :: acc_is_present_nb
+    module procedure :: acc_is_present_b
   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_copyout
+    module procedure :: acc_copyout_nb
+    module procedure :: acc_copyout_b
+  end interface
+  interface acc_copyout_async
+    module procedure :: acc_copyout_async_nb
+    module procedure :: acc_copyout_async_b
+  end interface
+  interface acc_copyout_finalize
+    module procedure :: acc_copyout_finalize_nb
+    module procedure :: acc_copyout_finalize_b
+  end interface
+  interface acc_copyout_finalize_async
+    module procedure :: acc_copyout_finalize_async_nb
+    module procedure :: acc_copyout_finalize_async_b
+  end interface
+
+  interface acc_delete
+    module procedure :: acc_delete_nb
+    module procedure :: acc_delete_b
+  end interface
+  interface acc_delete_async
+    module procedure :: acc_delete_async_nb
+    module procedure :: acc_delete_async_b
+  end interface
+  interface acc_delete_finalize
+    module procedure :: acc_delete_finalize_nb
+    module procedure :: acc_delete_finalize_b
+  end interface
+  interface acc_delete_finalize_async
+    module procedure :: acc_delete_finalize_async_nb
+    module procedure :: acc_delete_finalize_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
+contains
 
-  ! 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
+  subroutine acc_init(dev_type) 
+    use iso_c_binding
+    implicit none
+    !
+    integer(kind=acc_device_kind),value :: dev_type
+      interface
+        subroutine acc_init_c_impl(dev_type) &
+          bind(c,name="acc_init")
+          use iso_c_binding
+          Import::acc_device_kind
+          implicit none
+          !
+          integer(kind=acc_device_kind), value, intent(in):: dev_type
+        end subroutine
+      end interface
+      call acc_init_c_impl(dev_type)
+  end subroutine
 
-  ! Non-standard extensions
+  subroutine acc_shutdown(dev_type) 
+    use iso_c_binding
+    implicit none
+    !
+    integer(kind=acc_device_kind),value :: dev_type
+      interface
+        subroutine acc_shutdown_c_impl(dev_type) &
+          bind(c,name="acc_shutdown")
+          use iso_c_binding
+          Import::acc_device_kind
+          implicit none
+          !
+          integer(kind=acc_device_kind), value, intent(in):: dev_type
+        end subroutine
+      end interface
+      call acc_shutdown_c_impl(dev_type)
+  end subroutine
 
-  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
+  subroutine acc_set_device_type(dev_type) 
+      use iso_c_binding
       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
+      !
+      integer(kind=acc_device_kind),value :: dev_type
+      interface
+        subroutine acc_set_device_type_c_impl(dev_type) &
+          bind(c,name="acc_set_device_type")
+          use iso_c_binding
+          Import::acc_device_kind
+          implicit none
+          !
+          integer(kind=acc_device_kind), value, intent(in):: dev_type
+        end subroutine
+      end interface
+      call acc_set_device_type_c_impl(dev_type)
+  end subroutine
 
-contains
-  integer function acc_get_num_devices(dev_type)
+  integer(kind=acc_device_kind) function acc_get_device_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")
+    integer(c_int) function acc_get_device_type_c_impl() &
+        bind(c,name="acc_get_device_type")
+        use iso_c_binding
         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))
+    acc_get_device_type = int(acc_get_device_type_c_impl(),kind=acc_device_kind)
   end function
-    
-  subroutine acc_set_device_num(dev_num, dev_type)
+
+  subroutine acc_set_device_num(dev_num, dev_type) 
+    use iso_c_binding
     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")
+    !
+    integer(kind=acc_device_kind),value, intent(in) :: dev_type
+    integer,value, intent(in) :: dev_num
+    interface
+      subroutine acc_set_device_num_c_impl(dev_num, dev_type) &
+        bind(c,name="acc_set_device_num_f")
+        use iso_c_binding
+        Import::acc_device_kind
         implicit none
-         integer(c_int),value :: dev_num
-         integer(acc_device_kind),value :: dev_type
-      end subroutine         
+        !
+        integer(kind=acc_device_kind), value, intent(in):: dev_type
+        integer(c_int),value, intent(in) :: dev_num
+      end subroutine
     end interface
-    call acc_set_device_num_c_impl(int(dev_num,kind=c_int), dev_type)
-  end subroutine         
+    call acc_set_device_num_c_impl(dev_num, dev_type)
+  end subroutine
 
   integer function acc_get_device_num(dev_type)
+    use iso_c_binding
     implicit none
-    integer(acc_device_kind) :: dev_type
+    !
+    integer(kind=acc_device_kind),value :: dev_type
     interface
-      integer(c_int) function acc_get_device_num_c_impl(dev_type) bind(c,name="acc_get_device_num")
+      integer(c_int) function acc_get_device_num_c_impl(dev_type) &
+        bind(c,name="acc_get_device_num_f")
+        use iso_c_binding  
+        Import::acc_device_kind
         implicit none
+        !
+        integer(kind=acc_device_kind),value :: dev_type
       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)
+  integer(c_size_t) function acc_get_property(dev_num, dev_type, property) 
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: wait_arg
+    !
+    integer(kind=acc_device_kind),value, intent(in) :: dev_type
+    integer,value, intent(in) :: dev_num
+    integer(kind=acc_property_kind),value, intent(in) :: property
     interface
-      integer(c_int) function acc_async_test_c_impl(wait_arg) bind(c,name="acc_async_test")
+      integer(c_size_t) function acc_get_property_c_impl(dev_num, dev_type, property) &
+        bind(c,name="acc_get_property_f")
+        use iso_c_binding
+        Import::acc_device_kind, acc_property_kind
         implicit none
-        integer(acc_handle_kind),value :: wait_arg
+        !
+        integer(kind=acc_device_kind), value, intent(in):: dev_type
+        integer(c_int),value, intent(in) :: dev_num
+        integer(kind=acc_property_kind), value, intent(in):: property
       end function
     end interface
-    acc_async_test = acc_async_test_c_impl(wait_arg) > 0
+    acc_get_property = acc_get_property_c_impl(dev_num, dev_type, property)
   end function
-  
-  logical function acc_async_test_device(wait_arg, dev_num)
+
+  integer function acc_get_num_devices(dev_type)
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: wait_arg
-    integer :: dev_num
+    !
+    integer(kind=acc_device_kind),value :: dev_type
     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
+      integer(c_int) function acc_get_num_devices_c_impl(dev_type) bind(c,name="acc_get_num_devices")
+        use iso_c_binding  
+        Import::acc_device_kind
         implicit none
-        integer(acc_handle_kind),value :: wait_arg
-        integer(c_int),value :: dev_num
+        !
+        integer(kind=acc_device_kind),value :: dev_type
       end function
     end interface
-    acc_async_test_device = acc_async_test_device_c_impl(wait_arg,int(dev_num,kind=c_int)) > 0
+    acc_get_num_devices = int(acc_get_num_devices_c_impl(dev_type))
   end function
   
-  logical function acc_async_test_all()
+  logical function acc_is_present_b(data_arg, bytes)
+    use iso_c_binding
     implicit none
-    interface
-      integer(c_int) function acc_async_test_all_c_impl() &
-        bind(c,name=acc_async_test_all)
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      integer(c_int) function acc_is_present_c_impl(data_arg,bytes) &
+        bind(c,name="acc_is_present")
+        use iso_c_binding
         implicit none
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end function
     end interface
-    acc_async_test_all = acc_async_test_all_c_impl() > 0
+    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_async_test_all_device(dev_num)
+ 
+  logical function acc_is_present_nb(data_arg)
+    use iso_c_binding
     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
+    !
+    type(*), target, dimension(..), contiguous :: data_arg
+    acc_is_present_nb = acc_is_present_b(c_loc(data_arg),int(sizeof(data_arg)))
   end function
-  
-  subroutine acc_wait_device(wait_arg, dev_num)
+
+  subroutine acc_wait(wait_arg)
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: wait_arg
-    integer :: dev_num
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: wait_arg
+    !
     interface
-      subroutine acc_wait_device_c_impl(wait_arg, dev_num) &
-              bind(c,name="acc_wait_device")
-        use iso_c_binding, only: c_int
+      subroutine acc_wait_c_impl(wait_arg) &
+        bind(c,name="acc_wait")
+        use iso_c_binding
         implicit none
-        integer(acc_handle_kind),value :: wait_arg
-        integer(c_int),value :: dev_num
+        !
+        type(c_ptr),value,intent(in) :: wait_arg
+        !
       end subroutine
     end interface
-    call acc_wait_device_c_impl(wait_arg,int(dev_num,kind=c_int))
+    !
+    call acc_wait_c_impl(c_loc(wait_arg))
   end subroutine
-  
+
   subroutine acc_wait_async(wait_arg, async_arg)
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: wait_arg, async_arg
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: 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
+        bind(c,name="acc_wait_async")
+        use iso_c_binding
+        implicit none
+        !
+        type(c_ptr),value,intent(in) :: wait_arg, async_arg
+        !
       end subroutine
     end interface
-    call acc_wait_async_c_impl(wait_arg,async_arg)
+    !
+    call acc_wait_async_c_impl(c_loc(wait_arg), c_loc(async_arg))
   end subroutine
-  
+
+  subroutine acc_wait_all()
+    use iso_c_binding
+    implicit none
+    !
+    interface
+      subroutine acc_wait_all_c_impl() &
+        bind(c,name="acc_wait_all")
+        use iso_c_binding
+        implicit none
+        !
+      end subroutine
+    end interface
+    !
+    call acc_wait_all_c_impl()
+  end subroutine
+
+  subroutine acc_wait_all_async(async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    !
+    interface
+      subroutine acc_wait_all_async_c_impl(async_arg) &
+        bind(c,name="acc_wait_all_async")
+        use iso_c_binding
+        implicit none
+        !
+        type(c_ptr),value,intent(in) :: async_arg
+        !
+      end subroutine
+    end interface
+    !
+    call acc_wait_all_async_c_impl(c_loc(async_arg))
+  end subroutine
+
+  subroutine acc_wait_device(wait_arg, dev_num)
+    use iso_c_binding
+    implicit none
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: wait_arg
+    integer(c_int):: dev_num
+    !
+    interface
+      subroutine acc_wait_device_c_impl(wait_arg, dev_num) &
+        bind(c,name="acc_wait_device")
+        use iso_c_binding
+        implicit none
+        !
+        type(c_ptr),value,intent(in) :: wait_arg
+        integer(c_int),value,intent(in) :: dev_num
+        !
+      end subroutine
+    end interface
+    !
+    call acc_wait_device_c_impl(c_loc(wait_arg), dev_num)
+  end subroutine
+
   subroutine acc_wait_device_async(wait_arg, async_arg, dev_num)
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: wait_arg, async_arg
-    integer :: dev_num
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: wait_arg, async_arg
+    integer(c_int):: 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
+        bind(c,name="acc_wait_device_async")
+        use iso_c_binding
         implicit none
-        integer(acc_handle_kind),value :: wait_arg, async_arg
-        integer(c_int),value :: dev_num
+        !
+        type(c_ptr),value,intent(in) :: wait_arg, async_arg
+        integer(c_int),value,intent(in) :: dev_num
+        !
       end subroutine
     end interface
-    call acc_wait_device_async_c_impl(wait_arg,async_arg,int(dev_num,kind=c_int))
+    !
+    call acc_wait_device_async_c_impl(c_loc(wait_arg), c_loc(async_arg), dev_num)
   end subroutine
-  
+
   subroutine acc_wait_all_device(dev_num)
+    use iso_c_binding
     implicit none
-    integer :: dev_num
+    !
+    integer(c_int) :: 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
+        bind(c,name="acc_wait_all_device")
+        use iso_c_binding
         implicit none
-        integer(c_int),value :: dev_num
+        !
+        integer(c_int),value,intent(in) :: dev_num
+        !
       end subroutine
     end interface
-    call acc_wait_all_device_c_impl(int(dev_num,kind=c_int))
+    !
+    call acc_wait_all_device_c_impl(dev_num)
   end subroutine
-  
+
   subroutine acc_wait_all_device_async(async_arg, dev_num)
+    use iso_c_binding
     implicit none
-    integer(acc_handle_kind) :: async_arg
-    integer :: dev_num
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    integer(c_int):: 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
+        bind(c,name="acc_wait_all_device_async")
+        use iso_c_binding
         implicit none
-        integer(acc_handle_kind),value :: async_arg
-        integer(c_int),value :: dev_num
+        !
+        type(c_ptr),value,intent(in) :: async_arg
+        integer(c_int),value,intent(in) :: dev_num
+        !
       end subroutine
     end interface
-    call acc_wait_all_device_async_c_impl(async_arg,int(dev_num,kind=c_int))
+    !
+    call acc_wait_all_device_async_c_impl(c_loc(async_arg), dev_num)
   end subroutine
 
-  logical function acc_on_device(dev_type)
+  logical function acc_async_test(wait_arg)
+    use iso_c_binding
     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
+    !
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: wait_arg
+    !
+    interface
+      integer(c_int) function acc_async_test_c_impl(wait_arg) &
+        bind(c,name="acc_async_test")
+        use iso_c_binding
         implicit none
-        integer(acc_device_kind),value :: dev_type
+        !
+        type(c_ptr),value,intent(in) :: wait_arg
+        !
       end function
     end interface
-    acc_on_device = acc_on_device_c_impl(dev_type) > 0
+    !
+    acc_async_test = acc_async_test_c_impl(c_loc(wait_arg)) > 0
   end function
 
-  !subroutine acc_copyin(data_arg) # TODO
-  !subroutine acc_copyin_async(data_arg, async_arg) # TODO
-  subroutine acc_copyin(data_arg, bytes)
+  logical function acc_async_test_device(wait_arg, dev_num)
+    use iso_c_binding
     implicit none
-    type(*), target, dimension(..) :: data_arg
-    integer :: bytes
     !
-    type(c_ptr) :: tmp
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: wait_arg
+    integer(c_int) :: dev_num
+    !
     interface
-      function acc_copyin_c_impl(data_arg, bytes) bind(c,name="acc_copyin")
+      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
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        type(c_ptr) :: acc_copyin_c_impl
+        !
+        type(c_ptr),value,intent(in) :: wait_arg
+        integer(c_int) :: dev_num
+        !
       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)
+    !
+    acc_async_test_device = acc_async_test_device_c_impl(c_loc(wait_arg), dev_num) > 0
+  end function
+
+  logical function acc_async_test_all()
+    use iso_c_binding
+    implicit none
+    !
+    interface
+      integer(c_int) function acc_async_test_all_c_impl() &
+        bind(c,name="acc_async_test_all")
+        use iso_c_binding
+        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)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..) :: data_arg
-    integer :: bytes
-    integer(acc_handle_kind) :: async_arg
+    !
+    integer(c_int), value, intent(in) :: dev_num
+    !
     interface
-      subroutine acc_copyin_async_c_impl(data_arg, bytes) bind(c,name="acc_copyin_async")
+      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
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        integer(c_int), value, intent(in) :: dev_num
+        !
+      end function
+    end interface
+    !
+    acc_async_test_all_device = acc_async_test_all_device_c_impl(dev_num) > 0
+  end function
+
+  subroutine acc_copyin_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_copyin_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_copyin")
+        use iso_c_binding
+        implicit none
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end subroutine
     end interface
-    call acc_copyin_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyin_b_c_impl(c_loc(data_arg),int(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)
+
+  subroutine acc_copyin_nb(data_arg)
+    use iso_c_binding
     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")
+    type(*), target, dimension(..)::data_arg
+    call acc_copyin_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_copyin_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_copyin_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_copyin_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        type(c_ptr) :: acc_create_c_impl
-      end function
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
+      end subroutine
     end interface
-    tmp =  acc_create_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyin_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   end subroutine
-  subroutine acc_create_async(data_arg, bytes, async_arg)
+
+  subroutine acc_copyin_async_nb(data_arg, async_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_copyin_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  subroutine acc_copyout_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_copyout_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_copyout")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end subroutine
     end interface
-    call acc_create_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyout_b_c_impl(c_loc(data_arg),int(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)
+
+  subroutine acc_copyout_nb(data_arg)
+    use iso_c_binding
     implicit none
-    type(*), target, dimension(..) :: data_arg
-    integer :: bytes
-    interface
-      subroutine acc_copyout_c_impl(data_arg, bytes) bind(c,name="acc_copyout")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_copyout_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_copyout_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_copyout_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_copyout_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_copyout_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyout_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   end subroutine
-  subroutine acc_copyout_async(data_arg, bytes, async_arg)
+
+  subroutine acc_copyout_async_nb(data_arg, async_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_copyout_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  subroutine acc_copyout_finalize_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_copyout_finalize_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_copyout_finalize")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end subroutine
     end interface
-    call acc_copyout_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyout_finalize_b_c_impl(c_loc(data_arg),int(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)
+
+  subroutine acc_copyout_finalize_nb(data_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_copyout_finalize_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_copyout_finalize_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_copyout_finalize_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_copyout_finalize_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_copyout_finalize_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_copyout_finalize_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   end subroutine
-  subroutine acc_copyout_finalize_async(data_arg, bytes, async_arg)
+
+  subroutine acc_copyout_finalize_async_nb(data_arg, async_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..) :: data_arg
-    integer :: bytes
-    integer(acc_handle_kind) :: async_arg
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_copyout_finalize_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  integer(kind=acc_handle_kind) function acc_get_default_async() 
+    use iso_c_binding
+    implicit none
+    !
     interface
-      subroutine acc_copyout_finalize_async_c_impl(data_arg, bytes) bind(c,name="acc_copyout_finalize_async")
+    integer(c_int) function acc_get_default_async_c_impl() &
+        bind(c,name="acc_get_default_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+      end function
+    end interface
+    acc_get_default_async = int(acc_get_default_async_c_impl(),kind=acc_handle_kind)
+  end function
+
+  subroutine acc_set_default_async(async_arg) 
+    use iso_c_binding
+    implicit none
+    !
+    integer(kind=acc_handle_kind),value :: async_arg
+    interface
+      subroutine acc_set_default_async_c_impl(async_arg) &
+        bind(c,name="acc_set_default_async")
+        use iso_c_binding
+        Import::acc_handle_kind
+        implicit none
+        !
+        integer(kind=acc_handle_kind), value, intent(in):: async_arg
       end subroutine
     end interface
-    call acc_copyout_finalize_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_set_default_async_c_impl(async_arg)
   end subroutine
-  
-  !subroutine acc_delete(data_arg) # TODO
-  !subroutine acc_delete_async(data_arg, async_arg) # TODO
-  subroutine acc_delete(data_arg, bytes)
+
+  subroutine acc_create_b(data_arg, bytes)
+    use iso_c_binding
     implicit none
-    type(*), target, dimension(..) :: data_arg
-    integer :: bytes
-    interface
-      subroutine acc_delete_c_impl(data_arg, bytes) bind(c,name="acc_delete")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_create_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_create")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end subroutine
     end interface
-    call acc_delete_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_create_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
   end subroutine
-  subroutine acc_delete_async(data_arg, bytes, async_arg)
+
+  subroutine acc_create_nb(data_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_create_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_create_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_create_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_create_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_delete_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_create_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   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)
+
+  subroutine acc_create_async_nb(data_arg, async_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_create_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  subroutine acc_delete_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_delete_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_delete")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
       end subroutine
     end interface
-    call acc_delete_finalize_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_delete_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
   end subroutine
-  subroutine acc_delete_finalize_async(data_arg, bytes, async_arg)
+
+  subroutine acc_delete_nb(data_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_delete_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_delete_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_delete_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_delete_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: hostptr
-        integer(c_size_t),value :: num_bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_delete_finalize_async_c_impl(c_loc(data_arg), int(num_bytes,kind=c_size_t))
+    call acc_delete_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   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)
+  subroutine acc_delete_async_nb(data_arg, async_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..) :: data_arg
-    integer :: bytes
-    interface
-      subroutine acc_update_device_c_impl(data_arg, bytes) bind(c,name="acc_update_device")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_delete_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  subroutine acc_delete_finalize_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_delete_finalize_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_delete_finalize")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: data_arg
-        integer(c_size_t),value :: bytes
+        !
+        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))
+    call acc_delete_finalize_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
   end subroutine
 
-  subroutine acc_update_device_async(data_arg, bytes, async_arg)
+  subroutine acc_delete_finalize_nb(data_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_delete_finalize_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_delete_finalize_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_delete_finalize_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_delete_finalize_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: data_arg
-        integer(c_size_t),value :: bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_update_device_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t),async_arg)
+    call acc_delete_finalize_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   end subroutine
-  
-  subroutine acc_update_self_b(data_arg, bytes)
+
+  subroutine acc_delete_finalize_async_nb(data_arg, async_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..) :: data_arg
-    integer :: bytes
-    interface
-      subroutine acc_update_self_c_impl(data_arg, bytes) bind(c,name="acc_update_self")
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_delete_finalize_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
+  subroutine acc_update_device_b(data_arg, bytes)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_update_device_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_update_device")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: data_arg
-        integer(c_size_t),value :: bytes
+        !
+        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))
+    call acc_update_device_b_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)
+  subroutine acc_update_device_nb(data_arg)
+    use iso_c_binding
     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")
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_update_device_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_update_device_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    interface 
+      subroutine acc_update_device_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_update_device_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: data_arg
-        integer(c_size_t),value :: bytes
-        integer(acc_handle_kind),value :: async_arg
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
       end subroutine
     end interface
-    call acc_update_self_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t),async_arg)
+    call acc_update_device_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
   end subroutine
-  
-  subroutine acc_update_self_b(data_arg) 
+
+  subroutine acc_update_device_async_nb(data_arg, async_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..), contiguous :: data_arg
-    call acc_update_self_b(data_arg,sizeof(data_arg))
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_update_device_async_b(data_arg,int(sizeof(data_arg)), async_arg)
   end subroutine
-  subroutine acc_update_self_async_nb(data_arg, async_arg)
+
+  subroutine acc_update_self_b(data_arg, bytes)
+    use iso_c_binding
     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)
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    interface 
+      subroutine acc_update_self_b_c_impl(data_arg,bytes) &
+      bind(c,name="acc_update_self")
+        use iso_c_binding
+        implicit none
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+      end subroutine
+    end interface
+    call acc_update_self_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t))
   end subroutine
 
-  logical function acc_is_present_b(data_arg, bytes)
+  subroutine acc_update_self_nb(data_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..) :: data_arg
-    integer :: bytes
+    !
+    type(*), target, dimension(..)::data_arg
+    call acc_update_self_b(data_arg, int(sizeof(data_arg)))
+  end subroutine
+
+  subroutine acc_update_self_async_b(data_arg, bytes, async_arg)
+    use iso_c_binding
+    implicit none
+    !
+    type(*), target, dimension(..)::data_arg
+    integer,value,intent(in) :: bytes
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
     interface 
-      integer(c_int) function acc_is_present_c_impl(data_arg,bytes) bind(c,name="acc_is_present")
+      subroutine acc_update_self_async_b_c_impl(data_arg,bytes, async_arg) &
+      bind(c,name="acc_update_self_async")
+        use iso_c_binding
         implicit none
-        type(c_ptr),value :: data_arg
-        integer(c_size_t),value :: bytes
-      end function
+        !
+        type(c_ptr), value::data_arg
+        integer(c_size_t), value :: bytes
+        type(c_ptr),value,intent(in) :: async_arg
+      end subroutine
     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)
+    call acc_update_self_async_b_c_impl(c_loc(data_arg),int(bytes,kind=c_size_t), c_loc(async_arg))
+  end subroutine
+
+  subroutine acc_update_self_async_nb(data_arg, async_arg)
+    use iso_c_binding
     implicit none
-    type(*), dimension(..), contiguous :: data_arg
-    acc_is_present_nb = acc_is_present_b(data_arg,sizeof(data_arg)) 
-  end function
+    !
+    type(*), target, dimension(..)::data_arg
+    integer(acc_handle_kind),dimension(..),target,intent(in) :: async_arg
+    call acc_update_self_async_b(data_arg,int(sizeof(data_arg)), async_arg)
+  end subroutine
+
 end module
diff --git a/runtime/gpufortrt/test/cpp/test_acc_get_num_devices_00.cpp b/runtime/gpufortrt/test/cpp/test_acc_get_num_devices_00.cpp
new file mode 100755
index 00000000..0b7024b7
--- /dev/null
+++ b/runtime/gpufortrt/test/cpp/test_acc_get_num_devices_00.cpp
@@ -0,0 +1,10 @@
+#include <openacc.h>
+#include <iostream>
+
+int main( int argc, char * argv[] )
+{
+   acc_device_t t = acc_device_default;
+   std::cout<< "number of devices: " << acc_get_num_devices(t);
+
+   return 0;
+}
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/cpp/test_acc_init_shutdown_00.cpp b/runtime/gpufortrt/test/cpp/test_acc_init_shutdown_00.cpp
new file mode 100755
index 00000000..c8b90c0d
--- /dev/null
+++ b/runtime/gpufortrt/test/cpp/test_acc_init_shutdown_00.cpp
@@ -0,0 +1,21 @@
+#include <openacc.h>
+
+int main( int argc, char * argv[] )
+{
+   acc_init( acc_device_default );
+
+   float a[100];
+
+   #pragma acc data copyout(a[0:100])
+
+   #pragma acc parallel
+   #pragma acc loop
+   for( int i = 0; i < 100; ++i )
+   {
+      a[i] = 5;
+   }
+
+   acc_shutdown( acc_device_default );
+
+   return 0;
+}
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/cpp/test_acc_routine_00.cpp b/runtime/gpufortrt/test/cpp/test_acc_routine_00.cpp
new file mode 100755
index 00000000..de6e9772
--- /dev/null
+++ b/runtime/gpufortrt/test/cpp/test_acc_routine_00.cpp
@@ -0,0 +1,29 @@
+#include <openacc.h>
+
+#pragma acc routine worker
+void set( int* in_out )
+{
+   *in_out = ( *in_out ) * 3;
+}
+
+int main( int argc, char * argv[] )
+{
+   acc_init( acc_device_default );
+
+   float a[100];
+
+#pragma acc data copyout(a[0:100])
+
+#pragma acc parallel
+#pragma acc loop
+   for( int i = 0; i < 100; ++i )
+   {
+      int j = 5;
+      set(&j);
+      a[i] = j;
+   }
+
+   acc_shutdown( acc_device_default );
+
+   return 0;
+}
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/Makefile b/runtime/gpufortrt/test/fortran/Makefile
index 2af4ac8b..6cd11a03 100644
--- a/runtime/gpufortrt/test/fortran/Makefile
+++ b/runtime/gpufortrt/test/fortran/Makefile
@@ -10,6 +10,25 @@ LDFLAGS = $(shell gpufort --print-acc-ldflags) -lgfortran
 
 # sources
 TEST_SRC      = \
+				./test_acc_map_data_00.f90 \
+				./test_acc_malloc_00.f90 \
+				./test_acc_get_num_devices_00.f90 \
+				./test_acc_update_00.f90 \
+				./test_acc_delete_00.f90 \
+				./test_acc_create_00.f90 \
+				./test_acc_create_01.f90 \
+				./test_acc_set_get_default_async_00.f90 \
+				./test_acc_copyinout_00.f90 \
+				./test_acc_copyinout_01.f90 \
+				./test_acc_copyin_00.f90 \
+				./test_acc_async_test_00.f90 \
+				./test_acc_get_property_00.f90 \
+				./test_acc_is_present_00.f90 \
+				./test_acc_is_present_01.f90 \
+				./test_acc_get_device_num_00.f90 \
+				./test_acc_set_device_num_00.f90 \
+				./test_acc_init_shutdown_00.f90 \
+				./test_acc_wait_device_00.f90 \
 				./test_structured_region_00.f90 \
 				./test_structured_region_01.f90 \
 				./test_structured_region_02.f90 \
@@ -36,8 +55,8 @@ TEST_CONV_SRC = $(TEST_SRC:.f90=.f90-gpufort.f08)
 all: $(TEST_NAME)
 
 $(TEST_NAME): %: %.f90-gpufort.f08
-	$(HIPCC) $(HIPCC_CFLAGS) -c $@.f90-gpufort.cpp
-	$(HIPFC) $^ -o $@ $@.f90-gpufort.o $(HIPFC_CFLAGS) $(LDFLAGS)
+	$(HIPCC) $(HIPCC_CFLAGS) -c -g $@.f90-gpufort.cpp
+	$(HIPFC) -g $^ -o $@ $@.f90-gpufort.o $(HIPFC_CFLAGS) $(LDFLAGS)
 
 $(TEST_CONV_SRC): %-gpufort.f08: %
 	gpufort --wrap $^ --dest hipgpufort
diff --git a/runtime/gpufortrt/test/fortran/test_acc_async_test_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_async_test_00.f90
new file mode 100644
index 00000000..9d130946
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_async_test_00.f90
@@ -0,0 +1,23 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer(4) :: x(N), y(N)
+    logical:: xPresent = .FALSE., yPresent = .FALSE. 
+    call acc_init (acc_device_default)
+
+    !$ACC DATA COPY(X(1:N))
+    xPresent = acc_is_present(x,N)
+    yPresent = acc_is_present(y,N)
+    if ( xPresent .AND. .not. yPresent) then 
+        print *, "FINE!"
+    else
+        ERROR STOP "Arrays is not present"
+    end if
+    
+    !$acc end data
+    if (.NOT. acc_async_test (n) ) print *, "not sync!"
+    call acc_shutdown (acc_device_host)
+  end program
+  
diff --git a/runtime/gpufortrt/test/fortran/test_acc_copyin_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_copyin_00.f90
new file mode 100644
index 00000000..d11ebecb
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_copyin_00.f90
@@ -0,0 +1,25 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 100
+    integer :: i
+    integer :: x(N), y(N)
+    logical:: xPresent = .FALSE., yPresent = .FALSE. 
+  
+    call acc_init(acc_device_default)
+    do i = 1, N
+      y(i) = 3
+      x(i) = 2
+    end do
+    call acc_copyin(x, N)
+    xPresent = acc_is_present(x,N)
+    yPresent = acc_is_present(y,N)
+    if ( xPresent .AND. .not. yPresent) then 
+        print *, "FINE!"
+    else
+        ERROR STOP "Arrays is not present"
+    end if
+    call acc_shutdown(acc_device_default)
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_copyinout_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_copyinout_00.f90
new file mode 100644
index 00000000..10711b89
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_copyinout_00.f90
@@ -0,0 +1,38 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer :: i
+    integer(4) :: x(N), y(N), y_exact(N)
+
+    do i = 1, N
+      y_exact(i) = 3
+    end do
+  
+    call acc_copyin(x,N*4)
+    call acc_copyin(y,N*4)
+
+    !$acc parallel loop present(x,y)
+    !
+    do i = 1, N
+      x(i) = 1
+      y(i) = 2
+    end do
+    
+    !$acc parallel loop
+    do i = 1, N
+      y(i) = x(i) + y(i)
+    end do
+  
+    call acc_copyout(y,N*4)
+    
+    do i = 1, N
+      if ( y_exact(i) .ne.&
+              y(i) ) ERROR STOP "GPU and CPU result do not match"
+    end do
+  
+    print *, "PASSED"
+  
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_copyinout_01.f90 b/runtime/gpufortrt/test/fortran/test_acc_copyinout_01.f90
new file mode 100644
index 00000000..68c511c4
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_copyinout_01.f90
@@ -0,0 +1,38 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer :: i
+    integer(4) :: x(N), y(N), y_exact(N)
+
+    do i = 1, N
+      y_exact(i) = 3
+    end do
+  
+    call acc_copyin(x)
+    call acc_copyin(y)
+
+    !$acc parallel loop present(x,y)
+    !
+    do i = 1, N
+      x(i) = 1
+      y(i) = 2
+    end do
+    
+    !$acc parallel loop
+    do i = 1, N
+      y(i) = x(i) + y(i)
+    end do
+  
+    call acc_copyout(y)
+    
+    do i = 1, N
+      if ( y_exact(i) .ne.&
+              y(i) ) ERROR STOP "GPU and CPU result do not match"
+    end do
+  
+    print *, "PASSED"
+  
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_create_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_create_00.f90
new file mode 100644
index 00000000..eae0aa44
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_create_00.f90
@@ -0,0 +1,38 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer :: i
+    integer(4) :: x(N), y(N), y_exact(N)
+
+    do i = 1, N
+      y_exact(i) = 3
+    end do
+  
+    call acc_create(x,N*4)
+    call acc_create(y,N*4)
+
+    !$acc parallel loop present(x,y)
+    !
+    do i = 1, N
+      x(i) = 1
+      y(i) = 2
+    end do
+    
+    !$acc parallel loop
+    do i = 1, N
+      y(i) = x(i) + y(i)
+    end do
+  
+    call acc_copyout(y,N*4)
+    
+    do i = 1, N
+      if ( y_exact(i) .ne.&
+              y(i) ) ERROR STOP "GPU and CPU result do not match"
+    end do
+  
+    print *, "PASSED"
+  
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_create_01.f90 b/runtime/gpufortrt/test/fortran/test_acc_create_01.f90
new file mode 100644
index 00000000..78ba0386
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_create_01.f90
@@ -0,0 +1,38 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer :: i
+    integer(4) :: x(N), y(N), y_exact(N)
+
+    do i = 1, N
+      y_exact(i) = 3
+    end do
+  
+    call acc_create_async(x,N*4, acc_async_noval)
+    call acc_create_async(y,N*4, acc_async_noval)
+
+    !$acc parallel loop present(x,y)
+    !
+    do i = 1, N
+      x(i) = 1
+      y(i) = 2
+    end do
+    
+    !$acc parallel loop
+    do i = 1, N
+      y(i) = x(i) + y(i)
+    end do
+  
+    call acc_copyout(y,N*4)
+    
+    do i = 1, N
+      if ( y_exact(i) .ne.&
+              y(i) ) ERROR STOP "GPU and CPU result do not match"
+    end do
+  
+    print *, "PASSED"
+  
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_delete_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_delete_00.f90
new file mode 100644
index 00000000..449cfbcf
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_delete_00.f90
@@ -0,0 +1,41 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer :: i
+    integer(4) :: x(N), y(N), y_exact(N)
+
+    do i = 1, N
+      y_exact(i) = 3
+    end do
+  
+    call acc_create(x,N*4)
+    call acc_create(y,N*4)
+
+    !$acc parallel loop present(x,y)
+    !
+    do i = 1, N
+      x(i) = 1
+      y(i) = 2
+    end do
+    
+    !$acc parallel loop
+    do i = 1, N
+      y(i) = x(i) + y(i)
+    end do
+  
+    call acc_copyout(y,N*4)
+    
+    do i = 1, N
+      if ( y_exact(i) .ne.&
+              y(i) ) ERROR STOP "GPU and CPU result do not match"
+    end do
+
+    call acc_delete(x,N*4)
+    call acc_delete_async(y, N*4, acc_async_noval)
+    
+    print *, "PASSED"
+  
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_get_device_num_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_get_device_num_00.f90
new file mode 100755
index 00000000..1b04cc28
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_get_device_num_00.f90
@@ -0,0 +1,7 @@
+program main
+    use openacc
+    implicit none
+    integer(acc_device_kind) :: dev_type
+    dev_type = acc_device_default
+    print *, acc_get_device_num(dev_type)
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_get_device_type_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_get_device_type_00.f90
new file mode 100755
index 00000000..ab28f107
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_get_device_type_00.f90
@@ -0,0 +1,7 @@
+program main
+    use openacc
+    implicit none
+    integer(acc_device_kind) :: dev_type
+    dev_type = acc_get_device_type()
+    print *, dev_type
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_get_num_devices_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_get_num_devices_00.f90
new file mode 100755
index 00000000..71feeb25
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_get_num_devices_00.f90
@@ -0,0 +1,7 @@
+program main
+    use openacc
+    implicit none
+    integer:: num
+    num = acc_get_num_devices(acc_device_default)
+    print *, num
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_get_property_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_get_property_00.f90
new file mode 100644
index 00000000..2eb2dd2c
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_get_property_00.f90
@@ -0,0 +1,23 @@
+program main
+    use openacc
+    use iso_c_binding
+    implicit none
+    !
+    integer :: dev_type = acc_device_default
+    integer :: dev_num
+    integer :: dev
+    integer(c_size_t) :: val
+    dev_num = acc_get_num_devices(dev_type)
+  
+    do dev = 0, dev_num
+       print *, "Device ", dev
+  
+       val = acc_get_property (dev, dev_type, acc_property_memory)
+       print *, "    Total memory: ", val
+       if (val < 0) then
+          print *, "acc_property_memory should not be negative."
+          stop 1
+       end if
+    end do
+end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_init_shutdown_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_init_shutdown_00.f90
new file mode 100755
index 00000000..299e94fc
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_init_shutdown_00.f90
@@ -0,0 +1,5 @@
+program main
+    use openacc
+    call acc_init(acc_device_default)
+    call acc_shutdown(acc_device_default)
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_is_present_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_is_present_00.f90
new file mode 100755
index 00000000..383e2eb5
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_is_present_00.f90
@@ -0,0 +1,21 @@
+program main
+    ! begin of program
+    use openacc
+    implicit none
+    integer, parameter :: N = 1000
+    integer(4) :: x(N), y(N)
+    logical:: xPresent = .FALSE., yPresent = .FALSE. 
+  
+    !$ACC DATA COPY(X(1:N))
+    xPresent = acc_is_present(x,N)
+    yPresent = acc_is_present(y,N)
+    if ( xPresent .AND. .not. yPresent) then 
+        print *, "FINE!"
+    else
+        ERROR STOP "Arrays is not present"
+    end if
+    
+    !$acc end data
+  
+  end program
+  
diff --git a/runtime/gpufortrt/test/fortran/test_acc_is_present_01.f90 b/runtime/gpufortrt/test/fortran/test_acc_is_present_01.f90
new file mode 100644
index 00000000..6c22e31c
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_is_present_01.f90
@@ -0,0 +1,46 @@
+program main
+  ! begin of program
+  use openacc
+  implicit none
+  integer, parameter :: N = 1000
+  integer :: i
+  integer(4) :: x(N), y(N), y_exact(N)
+  logical:: xPresent = .FALSE., yPresent = .FALSE. 
+
+  do i = 1, N
+    y_exact(i) = 3
+  end do
+
+  call acc_copyin(x)
+  call acc_copyin(y)
+
+  xPresent = acc_is_present(x)
+  yPresent = acc_is_present(y, int(sizeof(y)))
+  if ( xPresent .AND. yPresent) then 
+      print *, "FINE!"
+  else
+      ERROR STOP "Arrays is not present"
+  end if
+
+  !$acc parallel loop 
+  !
+  do i = 1, N
+    x(i) = 1
+    y(i) = 2
+  end do
+  
+  !$acc parallel loop
+  do i = 1, N
+    y(i) = x(i) + y(i)
+  end do
+
+  call acc_copyout(y)
+  
+  do i = 1, N
+    if ( y_exact(i) .ne.&
+            y(i) ) ERROR STOP "GPU and CPU result do not match"
+  end do
+
+  print *, "PASSED"
+
+end program
diff --git a/runtime/gpufortrt/test/fortran/test_acc_malloc_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_malloc_00.f90
new file mode 100644
index 00000000..8b097d0f
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_malloc_00.f90
@@ -0,0 +1,51 @@
+program main
+  use iso_c_binding
+  implicit none
+  type (c_ptr) :: cptr
+  integer, parameter :: N = 1000
+  integer, pointer :: fptr(:)
+  integer :: test_array(N)
+  integer :: i = 0
+
+  interface
+    type (c_ptr) function acc_malloc (num_bytes) &
+      bind (C)
+      use iso_c_binding
+      integer (c_size_t), value :: num_bytes
+    end function
+  end interface
+
+  cptr = acc_malloc (N * sizeof (fptr(N)))
+  call c_f_pointer (cptr, fptr, [N])
+    
+  call dummy_subroutine (fptr, test_array, N)
+    
+  do i = 1, N
+      if (test_array(i) .ne. i) ERROR STOP "Results do not match"
+  end do
+    
+  print *, "PASSED"
+
+  contains
+  subroutine dummy_subroutine (fptr, test_array, N)
+      use openacc
+      implicit none
+      integer :: N
+      integer :: test_array(N)
+      integer :: fptr(N)
+      integer :: i = 0
+      
+      !$acc data deviceptr (fptr)
+  
+      call acc_copyin(test_array)
+      !$acc parallel 
+        do i = 1, N
+          fptr(i) = i
+          test_array(i) = fptr(i)
+        end do
+      !$acc end parallel
+      call acc_copyout(test_array)
+      
+    end subroutine
+
+end program main  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_map_data_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_map_data_00.f90
new file mode 100644
index 00000000..75059c26
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_map_data_00.f90
@@ -0,0 +1,61 @@
+program main
+    use iso_c_binding
+    implicit none
+    type (c_ptr), target :: cptr
+    integer, parameter :: N = 5
+    integer, pointer :: fptr(:)
+    integer, target :: data_arg(N)
+    integer :: i = 0
+
+    interface
+    function acc_malloc (num_bytes) result(res) &
+      bind (C)
+      use iso_c_binding
+      integer (c_size_t), value :: num_bytes
+      !
+      type(c_ptr)  :: res
+    end function
+    end interface
+
+    interface
+    subroutine acc_map_data (data_arg, data_dev, num_bytes) &
+      bind (C)
+      use iso_c_binding
+      type (c_ptr), value:: data_arg
+      type (c_ptr), value:: data_dev
+      integer (c_size_t), value :: num_bytes
+    end subroutine
+    end interface
+
+    cptr = acc_malloc (N * sizeof (fptr(N)))
+    call c_f_pointer (cptr, fptr, [N])
+
+    call initialize_device_memory (fptr, N)
+
+    call acc_map_data (c_loc(data_arg), cptr, N * sizeof (fptr(N)))
+    
+    do i = 1, N
+        if (data_arg(i) .ne. i * 2) ERROR STOP "Results do not match"
+    end do
+      
+    print *, "PASSED"
+
+    contains
+    subroutine initialize_device_memory (fptr, N)
+      use openacc
+      implicit none
+      integer :: N
+      integer :: fptr(N)
+      integer :: i = 0
+      
+      !$acc data deviceptr (fptr)
+
+      !$acc parallel 
+        do i = 1, N
+          fptr(i) = i * 2
+        end do
+      !$acc end parallel
+
+    end subroutine
+  
+  end program main  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_set_device_num_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_set_device_num_00.f90
new file mode 100755
index 00000000..fb4b336b
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_set_device_num_00.f90
@@ -0,0 +1,6 @@
+program main
+    use openacc
+    implicit none
+    integer :: dev_num = 1
+    call acc_set_device_num(dev_num, acc_device_default)
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_set_device_type_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_set_device_type_00.f90
new file mode 100755
index 00000000..64dbe351
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_set_device_type_00.f90
@@ -0,0 +1,5 @@
+program main
+    use openacc
+    call acc_set_device_type(acc_device_default)
+    print *, acc_get_device_type()
+end program
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_set_get_default_async_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_set_get_default_async_00.f90
new file mode 100644
index 00000000..1a7296de
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_set_get_default_async_00.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+program main
+    use openacc
+    implicit none
+    integer :: i
+    integer, parameter :: N = 1000
+    integer(4) :: x(N)
+      
+    call acc_set_default_async(1)
+    call acc_copyin_async(x,N*4,acc_async_noval)
+
+    !$acc parallel  
+      do i = 1, N
+        x(i) = 1
+      end do
+    !$acc end parallel
+  
+    call acc_wait_async (0, acc_async_noval)
+  
+    ! Test unseen async-argument.
+    if (acc_async_test (12) .neqv. .TRUE.) stop 1
+    call acc_wait_async (12, acc_async_noval)
+  
+    call acc_wait (1)
+  
+    if (acc_async_test (0) .neqv. .TRUE.) stop 2
+    if (acc_async_test (1) .neqv. .TRUE.) stop 3
+    if (acc_async_test (2) .neqv. .TRUE.) stop 4
+  
+    print *, "acc_get_default_async: ", acc_get_default_async()
+    
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_update_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_update_00.f90
new file mode 100644
index 00000000..88d482d1
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_update_00.f90
@@ -0,0 +1,56 @@
+! See also "lib-16-2.f90".
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program main
+    use openacc
+    implicit none
+  
+    integer, parameter :: N = 1000
+    integer(4) :: x(N)
+    integer :: i
+    ! integer :: async = 5
+
+    call acc_set_default_async(5)
+    
+    do i = 1, N
+      x(i) = 3
+    end do 
+  
+    call acc_copyin (x)
+  
+    do i = 1, N
+      x(i) = i + 1
+    end do 
+  
+    call acc_update_device_async (x, 4*N, acc_async_noval)
+    
+    ! We must wait for the update to be done.
+    call acc_wait (acc_async_noval)
+    
+    call acc_copyout_async (x, 4*N, acc_async_noval)
+  
+    call acc_wait (acc_async_noval)
+  
+    do i = 1, N
+      if (x(i) /= i + 1) stop 2
+    end do 
+  
+    call acc_copyin (x, 4*N)
+    
+    call acc_update_self_async (x, 4*N, acc_async_noval)
+      
+    call acc_wait (acc_async_noval)
+  
+    do i = 1, N
+      if (x(i) /= i + 1) stop 4
+    end do 
+  
+    call acc_delete_async (x, acc_async_noval)
+  
+    call acc_wait (acc_async_noval)
+
+    print *, "PASSED"
+      
+  end program
+  
\ No newline at end of file
diff --git a/runtime/gpufortrt/test/fortran/test_acc_wait_device_00.f90 b/runtime/gpufortrt/test/fortran/test_acc_wait_device_00.f90
new file mode 100644
index 00000000..029d4b46
--- /dev/null
+++ b/runtime/gpufortrt/test/fortran/test_acc_wait_device_00.f90
@@ -0,0 +1,6 @@
+program main
+    use openacc
+    implicit none
+    integer :: n = 0
+    call acc_wait_device(n , 0)
+end program
\ No newline at end of file