merge aomp-dev into amd-staging

test/smoke-fort-dev/use-omp-get-mapped-ptr/Makefile

@@ -0,0 +1,23 @@
+include ../../Makefile.defs
+
+TESTNAME     = use_omp_get_mapped_ptr
+TESTSRC_MAIN = main.f90
+TESTSRC_AUX  = bar.o
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+AOMPHIP ?= $(AOMP)
+HIPCC ?= $(AOMPHIP)/bin/hipcc
+HIP_CLANG_PATH ?= $(AOMP)/bin
+
+CFLAGS = -O3
+FLANG        ?= flang
+OMP_BIN      = $(AOMP)/bin/$(FLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+EXTRA_CFLAGS = -L$(AOMPHIP)/lib -lamdhip64 -Wl,-rpath,$(AOMPHIP)/lib -fPIC
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
+all: $(TESTNAME)
+
+bar.o : bar.hip
+	HIP_CLANG_PATH=$(HIP_CLANG_PATH) $(HIPCC) -c --offload-arch=$(AOMP_GPU) -fPIC $^ -o $@

test/smoke-fort-dev/use-omp-get-mapped-ptr/bar.hip

@@ -0,0 +1,23 @@
+#include <hip/hip_runtime.h>
+#include <stdlib.h>
+
+__global__ void bar_kernel(int *x, int *y, int *z, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (i < n) {
+          z[i] += x[i] + y[i];
+    }
+}
+
+extern "C" {
+
+void bar_GPU(int *x, int *y, int *z, int n)
+{
+    int num_threads = 256;
+    int num_blocks = n / num_threads + 1;
+    hipLaunchKernelGGL(bar_kernel, dim3(num_blocks), dim3(num_threads), 0, 0, x, y, z, n);
+    (void)hipDeviceSynchronize();
+}
+
+} /* extern "C" */

test/smoke-fort-dev/use-omp-get-mapped-ptr/main.f90

@@ -0,0 +1,79 @@
+MODULE foo
+  USE iso_c_binding
+  USE omp_lib
+  IMPLICIT NONE
+  PRIVATE
+  PUBLIC :: bar_device_ptr, bar_device_addr
+
+  INTERFACE
+    SUBROUTINE bar(x, y, z, n) BIND(C, name="bar_GPU")
+      USE iso_c_binding
+      TYPE(C_PTR),    VALUE, INTENT(IN)    :: x, y, z
+      INTEGER(C_INT), VALUE, INTENT(IN)    :: n
+    END SUBROUTINE
+
+  END INTERFACE
+
+CONTAINS
+
+  SUBROUTINE bar_device_addr(x,y,z,n)
+    INTEGER, TARGET, INTENT(IN)    :: x(:), y(:)
+    INTEGER, TARGET, INTENT(INOUT) :: z(:)
+    INTEGER(C_INT), INTENT(IN)     :: n
+    !$omp target data use_device_addr (x, y, z)
+    CALL bar(c_loc(x), c_loc(y), c_loc(z), n)
+    !$omp end target data
+  END SUBROUTINE
+
+  SUBROUTINE bar_device_ptr(x,y,z,n)
+    INTEGER, TARGET, INTENT(IN)    :: x(:), y(:)
+    INTEGER, TARGET, INTENT(INOUT) :: z(:)
+    INTEGER(C_INT), INTENT(IN)     :: n
+    TYPE(C_PTR)                    :: x_ptr, y_ptr, z_ptr
+
+    x_ptr = omp_get_mapped_ptr(c_loc(x),  omp_get_default_device())
+    y_ptr = omp_get_mapped_ptr(c_loc(y),  omp_get_default_device())
+    z_ptr = omp_get_mapped_ptr(c_loc(z),  omp_get_default_device())
+    CALL bar(x_ptr, y_ptr, z_ptr, n)
+  END SUBROUTINE
+
+END MODULE foo
+
+PROGRAM test_ptr
+  USE iso_c_binding
+  USE omp_lib
+  USE foo
+  IMPLICIT NONE
+
+  INTEGER, ALLOCATABLE, TARGET :: x(:), y(:), z(:)
+  INTEGER, ALLOCATABLE, TARGET :: x1(:), y1(:), z1(:)
+  INTEGER(C_INT) :: i, n
+  n = 1000
+  ALLOCATE(x(n), y(n), z(n))
+  ALLOCATE(x1(n), y1(n), z1(n))
+  z = 0
+  z1 = 0
+  x = 1
+  y = 2
+  x1 = 1
+  y1 = 2
+  i = 1
+  !$omp target enter data map(to: x,y,z,x1,y1,z1)
+
+  CALL bar_device_addr(x,y,z,n)
+  CALL bar_device_ptr(x1,y1,z1,n)
+  !$omp target exit data map(from: x,y,z,x1,y1,z1)
+  DO i = 1,n
+     IF (z(i) .ne. 3) then
+       PRINT *, "Bad result for use_device_addr!"
+       STOP 1
+     ENDIF
+     IF (z1(i) .ne. 3) then
+       PRINT *, "Bad result for omp_get_mapped_ptr!"
+       STOP 1
+     ENDIF
+  END DO
+  DEALLOCATE(x,y,z,x1,y1,z1)
+  PRINT *, "Success"
+END PROGRAM test_ptr
+

test/smoke-fort-fails/use-device-ptr-cptr/Makefile

@@ -0,0 +1,23 @@
+include ../../Makefile.defs
+
+TESTNAME     = use_device_ptr_cptr
+TESTSRC_MAIN = main.f90
+TESTSRC_AUX  = bar.o
+TESTSRC_ALL  = $(TESTSRC_MAIN) $(TESTSRC_AUX)
+AOMPHIP ?= $(AOMP)
+HIPCC ?= $(AOMPHIP)/bin/hipcc
+HIP_CLANG_PATH ?= $(AOMP)/bin
+
+CFLAGS = -O3
+FLANG        ?= flang
+OMP_BIN      = $(AOMP)/bin/$(FLANG)
+CC           = $(OMP_BIN) $(VERBOSE)
+EXTRA_CFLAGS = -L$(AOMPHIP)/lib -lamdhip64 -Wl,-rpath,$(AOMPHIP)/lib -fPIC
+#-ccc-print-phases
+#"-\#\#\#"
+
+include ../Makefile.rules
+all: $(TESTNAME)
+
+bar.o : bar.hip
+	HIP_CLANG_PATH=$(HIP_CLANG_PATH) $(HIPCC) -c --offload-arch=$(AOMP_GPU) -fPIC $^ -o $@

test/smoke-fort-fails/use-device-ptr-cptr/bar.hip

@@ -0,0 +1,23 @@
+#include <hip/hip_runtime.h>
+#include <stdlib.h>
+
+__global__ void bar_kernel(int *x, int *y, int *z, int n)
+{
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+    if (i < n) {
+          z[i] += x[i] + y[i];
+    }
+}
+
+extern "C" {
+
+void bar_GPU(int *x, int *y, int *z, int n)
+{
+    int num_threads = 256;
+    int num_blocks = n / num_threads + 1;
+    hipLaunchKernelGGL(bar_kernel, dim3(num_blocks), dim3(num_threads), 0, 0, x, y, z, n);
+    (void)hipDeviceSynchronize();
+}
+
+} /* extern "C" */

test/smoke-fort-fails/use-device-ptr-cptr/main.f90

@@ -0,0 +1,81 @@
+MODULE foo
+  USE iso_c_binding
+  USE omp_lib
+  IMPLICIT NONE
+  PRIVATE
+  PUBLIC :: bar_device_ptr, bar_device_addr
+
+  INTERFACE
+    SUBROUTINE bar(x, y, z, n) BIND(C, name="bar_GPU")
+      USE iso_c_binding
+      TYPE(C_PTR),    VALUE, INTENT(IN)    :: x, y, z
+      INTEGER(C_INT), VALUE, INTENT(IN)    :: n
+    END SUBROUTINE
+
+  END INTERFACE
+
+CONTAINS
+
+  SUBROUTINE bar_device_addr(x,y,z,n)
+    INTEGER, TARGET, INTENT(IN)    :: x(:), y(:)
+    INTEGER, TARGET, INTENT(INOUT) :: z(:)
+    INTEGER(C_INT), INTENT(IN)     :: n
+    !$omp target data use_device_addr (x, y, z)
+    CALL bar(c_loc(x), c_loc(y), c_loc(z), n)
+    !$omp end target data
+  END SUBROUTINE
+
+  SUBROUTINE bar_device_ptr(x,y,z,n)
+    INTEGER, TARGET, INTENT(IN)    :: x(:), y(:)
+    INTEGER, TARGET, INTENT(INOUT) :: z(:)
+    INTEGER(C_INT), INTENT(IN)     :: n
+    TYPE(C_PTR)                    :: x_ptr, y_ptr, z_ptr
+
+    x_ptr = c_loc(x)
+    y_ptr = c_loc(y)
+    z_ptr = c_loc(z)
+    !$omp target data use_device_ptr (x_ptr, y_ptr, z_ptr)
+    CALL bar(x_ptr, y_ptr, z_ptr, n)
+    !$omp end target data
+  END SUBROUTINE
+
+END MODULE foo
+
+PROGRAM test_ptr
+  USE iso_c_binding
+  USE omp_lib
+  USE foo
+  IMPLICIT NONE
+
+  INTEGER, ALLOCATABLE, TARGET :: x(:), y(:), z(:)
+  INTEGER, ALLOCATABLE, TARGET :: x1(:), y1(:), z1(:)
+  INTEGER(C_INT) :: i, n
+  n = 1000
+  ALLOCATE(x(n), y(n), z(n))
+  ALLOCATE(x1(n), y1(n), z1(n))
+  z = 0
+  z1 = 0
+  x = 1
+  y = 2
+  x1 = 1
+  y1 = 2
+  i = 1
+  !$omp target enter data map(to: x,y,z,x1,y1,z1)
+
+  CALL bar_device_addr(x,y,z,n)
+  CALL bar_device_ptr(x1,y1,z1,n)
+  !$omp target exit data map(from: x,y,z,x1,y1,z1)
+  DO i = 1,n
+     IF (z(i) .ne. 3) then
+       PRINT *, "Bad result for use_device_addr!"
+       STOP 1
+     ENDIF
+     IF (z1(i) .ne. 3) then
+       PRINT *, "Bad result for use_device_ptr!"
+       STOP 1
+     ENDIF
+  END DO
+  DEALLOCATE(x,y,z,x1,y1,z1)
+  PRINT *, "Success"
+END PROGRAM test_ptr
+