WiP: Create OpenACC runtime layer on top of GPUFORT RT

runtime/README.md

@@ -29,11 +29,17 @@
   * No thread-safe access (unless all required arrays are present)
   * No standard-compliant runtime API (might change long term) 
 
+  Requirements:
+  * `python3` with `jinja2` package is required to generate some of the interfaces
+     this runtime is using.
+  * Non-standard GNU Fortran extension `sizeof` must be supported by the used Fortran compiler
+  * Assumed-type, assumed-rank arrays must be supported by the used Fortran compiler
+
+
 ## Runtime subfolders
 
 * `openacc_gomp`: contains extended interfaces to GCC's LIBGOMP.
-* `gpufortrt`: contains a minimal non-standard-compliant runtime written 
-  completely in Fortran, which can be used teaching purposes
+* `gpufortrt`: contains a minimal non-standard-compliant runtime.
 
 ## Building
 
@@ -43,22 +49,6 @@ make
 ```
 
 Take a look into the `Makefile` in each folder and adapt it to your needs.
-Certain Makefile variables can be overwritten via the environment variables  
-`FC`, `HIPFORT_INC`, `FCFLAGS`, `CXX`, `CXXFLAGS`.
-
-Example:
-
-```
-cd <runtime_subfolder>
-FC=/usr/bin/gfortran HIPFORT_INC=/opt/rocm/hipfort/install/include/ make clean all
-```
-
-> **NOTE:** `python3` with `jinja2` package installed is required to use the codegen target.`
+Certain Makefile variables can be overwritten via environment variables.
 
-## Outlook
 
-* We require a standard-compliant, thread-safe HIP C++ based runtime that can
-  be used/built for HIP devices from different vendors (AMD, NVIDIA) to
-  support efforts such as GPUFORT. Main purpose of the runtime must be
-  mapping data between host and device and managing HIP streams.
-  Other features are optional.

runtime/gpufortrt/Makefile

@@ -2,7 +2,7 @@
 # Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
 include rules.mk
 
-CXXFLAGS += -I./src -I./src/internal
+CXXFLAGS += -I./src -I./src/internal -I./include
 
 VPATH = ./src:./src/internal
 
@@ -36,7 +36,7 @@ codegen:
 
 clean_all:
 	rm -f *.o *.mod *.a
-	rm -rf include/ lib/
+	rm -rf include/*.mod lib/
 
 clean:
 	rm -f *.o *.mod *.a

runtime/gpufortrt/src/gpufortrt_api.h → runtime/gpufortrt/include/gpufortrt_api.h

@@ -5,10 +5,22 @@
 #ifdef __cplusplus
 extern "C" {
 #endif
-  extern int gpufortrt_async_noval;
-
+  int gpufortrt_get_num_devices();
+
+  void gpufortrt_set_device_num(int dev_num);
+  int gpufortrt_get_device_num();
+
+  size_t gpufortrt_get_property(int dev_num,
+                                gpufortrt_device_property_t property);
+  const 
+  char* gpufortrt_get_property_string(int dev_num,
+                                      gpufortrt_device_property_t property);
+
   void gpufortrt_init();
   void gpufortrt_shutdown();
+
+  int gpufortrt_get_default_async(void);
+  void gpufortrt_set_default_async(int async_arg);
 
   void gpufortrt_data_start(
           gpufortrt_mapping_t* mappings,
@@ -77,7 +89,7 @@ extern "C" {
           void* hostptr,
           std::size_t num_bytes,
           bool never_deallocate);
-  void* gpufortrt_create_async(
+  void gpufortrt_create_async(
           void* hostptr,
           std::size_t num_bytes,
           int async_arg,
@@ -87,21 +99,11 @@ extern "C" {
           void* hostptr,
           std::size_t num_bytes,
           bool never_deallocate);
-  void* gpufortrt_copyin_async(
+  void gpufortrt_copyin_async(
           void* hostptr,
           std::size_t num_bytes,
           int async_arg,
           bool never_deallocate);
-
-  void* gpufortrt_copy(
-         void* hostptr,
-         std::size_t num_bytes,
-         bool never_deallocate);
-  void* gpufortrt_copy_async(
-         void* hostptr,
-         std::size_t num_bytes,
-         int async_arg,
-         bool never_deallocate);
 
   // other runtime calls
 
@@ -138,6 +140,11 @@ extern "C" {
                             bool if_arg);
   void gpufortrt_wait_all_async(int* async_arg,int num_async,
                                 bool if_arg);
+
+  int gpufortrt_async_test(int wait_arg);
+  int gpufortrt_async_test_device(int wait_arg, int dev_num);
+  int gpufortrt_async_test_all(void);
+  int gpufortrt_async_test_all_device(int dev_num);
 
   gpufortrt_queue_t gpufortrt_get_stream(int async_arg);
 

runtime/gpufortrt/src/gpufortrt_types.h → runtime/gpufortrt/include/gpufortrt_types.h

@@ -8,6 +8,19 @@
 #ifdef __cplusplus
 extern "C" {
 #endif
+  extern int gpufortrt_async_noval;
+  extern int gpufortrt_async_sync;
+  extern int gpufortrt_async_default;
+
+  enum gpufortrt_property_t {
+    gpufortrt_property_memory = 0,//>integer,  size of device memory in bytes
+    gpufortrt_property_free_memory,//>integer,  free device memory in bytes
+    gpufortrt_property_shared_memory_support,//>integer,  nonzero if the specified device supports sharing memory with the local thread
+    gpufortrt_property_name,//>string, device name*/
+    gpufortrt_property_vendor,//>string, device vendor*/
+    gpufortrt_property_driver,//>string, device driver version*/
+  };
+
   /**
    * \note: Enum values must match those of Fortran enumeration!
    * \note: Upper case first letter used because `delete` is C++ keyword.
@@ -56,4 +69,4 @@ std::ostream& operator<<(std::ostream& os, gpufortrt_map_kind_t map_kind);
 std::ostream& operator<<(std::ostream& os, gpufortrt_counter_t counter);
 std::ostream& operator<<(std::ostream& os,const gpufortrt_mapping_t& mapping);
 #endif
-#endif // GPUFORTRT_TYPES_H
+#endif // GPUFORTRT_TYPES_H

runtime/gpufortrt/include/openacc.h

@@ -0,0 +1,167 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
+#ifndef OPENACC_LIB_H
+#define OPENACC_LIB_H
+#ifdef __cplusplus
+extern "C" {
+#endif
+#include "gpufortrt_types.h"
+
+extern int acc_async_noval;
+extern int acc_async_sync;
+extern int acc_async_default;
+
+#define h_void void
+#define d_void void
+
+/** \note Enum values assigned according to `acc_set_device_num` description.*/
+enum acc_device_t {
+  acc_device_default = -1,
+  acc_device_all = 0,     
+  acc_device_none = 1,
+  acc_device_host,
+  acc_device_current,
+  acc_device_not_host,
+  acc_device_hip = acc_device_not_host,
+  acc_device_radeon = acc_device_hip,
+  acc_device_nvidia = acc_device_hip
+};
+
+enum acc_property_t {
+  acc_property_memory = 0,//>integer,  size of device memory in bytes
+  acc_property_free_memory,//>integer,  free device memory in bytes
+  acc_property_shared_memory_support,//>integer,  nonzero if the specified device supports sharing memory with the local thread
+  acc_property_name,//>string, device name*/
+  acc_property_vendor,//>string, device vendor*/
+  acc_property_driver//>string, device driver version*/
+};
+
+const char* ACC_DEVICE_TYPE_ENV_VAR = "ACC_DEVICE_TYPE"; //> Upper case identifer, maps to `acc_device_t` suffix.
+const char* ACC_DEVICE_NUM_ENV_VAR = "ACC_DEVICE_NUM"; //> Non-negative integer, number of default devices.
+//const char* ACC_PROFLIB_ENV_VAR = "ACC_PROFLIB";
+
+#define h_void void
+#define d_void void
+
+// except for acc_on_device all routines are only available on the host
+
+int acc_get_num_devices(acc_device_t dev_type);
+
+void acc_set_device_type(acc_device_t dev_type);
+acc_device_t acc_get_device_type(void);
+
+void acc_set_device_num(int dev_num, acc_device_t dev_type);
+int acc_get_device_num(acc_device_t dev_type);
+
+size_t acc_get_property(int dev_num,
+                        acc_device_t dev_type,
+                        acc_device_property_t property);
+const 
+char* acc_get_property_string(int dev_num,
+                              acc_device_t dev_type,
+                              acc_device_property_t property);
+
+void acc_init(acc_on_device_t dev_type);
+void acc_shutdown(acc_device_t dev_type);
+
+int acc_async_test(int wait_arg);
+int acc_async_test_device(int wait_arg, int dev_num);
+int acc_async_test_all(void);
+int acc_async_test_all_device(int dev_num);
+
+void acc_wait(int wait_arg);
+void acc_wait_device(int wait_arg, int dev_num);
+void acc_wait_async(int wait_arg, int async_arg);
+void acc_wait_device_async(int wait_arg, int async_arg,
+int dev_num);
+void acc_wait_all(void);
+void acc_wait_all_device(int dev_num);
+void acc_wait_all_async(int async_arg);
+void acc_wait_all_device_async(int async_arg, int dev_num);
+
+int acc_get_default_async(void);
+void acc_set_default_async(int async_arg);
+
+#ifdef __HIPCC__
+__host__ __device__ int acc_on_device(acc_device_t dev_type);
+#else
+int acc_on_device(acc_device_t dev_type);
+#endif
+
+d_void* acc_malloc(size_t bytes); // TODO create gpufortrt equivalent
+void acc_free(d_void* data_dev);  // TODO create gpufortr equivalent
+
+d_void* acc_copyin(h_void* data_arg, size_t bytes);
+void acc_copyin_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+
+d_void* acc_create(h_void* data_arg, size_t bytes);
+void acc_create_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+
+void acc_copyout(h_void* data_arg, size_t bytes);
+void acc_copyout_async(h_void* data_arg, size_t bytes,
+                       int async_arg);
+void acc_copyout_finalize(h_void* data_arg, size_t bytes);
+void acc_copyout_finalize_async(h_void* data_arg, size_t bytes,
+                                int async_arg);
+
+void acc_delete(h_void* data_arg, size_t bytes);
+void acc_delete_async(h_void* data_arg, size_t bytes,
+                      int async_arg);
+void acc_delete_finalize(h_void* data_arg, size_t bytes);
+void acc_delete_finalize_async(h_void* data_arg,
+                               size_t bytes, int async_arg);
+
+void acc_update_device(h_void* data_arg, size_t bytes);
+void acc_update_device_async(h_void* data_arg, size_t bytes,
+                             int async_arg);
+
+void acc_update_self(h_void* data_arg, size_t bytes);
+void acc_update_self_async(h_void* data_arg, size_t bytes,
+                           int async_arg);
+
+void acc_map_data(h_void* data_arg, d_void* data_dev,
+                  size_t bytes);
+void acc_unmap_data(h_void* data_arg);
+
+d_void* acc_deviceptr(h_void* data_arg);
+
+h_void* acc_hostptr(d_void* data_dev);
+
+int acc_is_present(h_void* data_arg, size_t bytes);
+
+void acc_memcpy_to_device(d_void* data_dev_dest,
+                          h_void* data_host_src, size_t bytes);
+void acc_memcpy_to_device_async(d_void* data_dev_dest,
+                                h_void* data_host_src, size_t bytes,
+                                int async_arg);
+
+void acc_memcpy_from_device(h_void* data_host_dest,
+                            d_void* data_dev_src, size_t bytes);
+void acc_memcpy_from_device_async(h_void* data_host_dest,
+                                  d_void* data_dev_src, size_t bytes,
+                                  int async_arg)
+
+void acc_attach(h_void** ptr_addr);
+void acc_attach_async(h_void** ptr_addr, int async_arg);
+
+void acc_detach(h_void** ptr_addr);
+void acc_detach_async(h_void** ptr_addr, int async_arg);
+void acc_detach_finalize(h_void** ptr_addr);
+void acc_detach_finalize_async(h_void** ptr_addr,
+                               int async_arg);
+
+void acc_memcpy_d2d(h_void* data_arg_dest,
+                    h_void* data_arg_src, size_t bytes,
+                    int dev_num_dest, int dev_num_src);
+
+void acc_memcpy_d2d_async(h_void* data_arg_dest,
+                          h_void* data_arg_src, size_t bytes,
+                          int dev_num_dest, int dev_num_src,
+                          int async_arg_src);
+
+#ifdef __cplusplus
+} // extern "C"
+#endif
+#endif // OPENACC_LIB_H