diff --git a/test/smoke-fort-dev/use-omp-get-mapped-ptr/Makefile b/test/smoke-fort-dev/use-omp-get-mapped-ptr/Makefile new file mode 100644 index 000000000..8f8af71d4 --- /dev/null +++ b/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 $@ diff --git a/test/smoke-fort-dev/use-omp-get-mapped-ptr/bar.hip b/test/smoke-fort-dev/use-omp-get-mapped-ptr/bar.hip new file mode 100644 index 000000000..659aa8828 --- /dev/null +++ b/test/smoke-fort-dev/use-omp-get-mapped-ptr/bar.hip @@ -0,0 +1,23 @@ +#include +#include + +__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" */ diff --git a/test/smoke-fort-dev/use-omp-get-mapped-ptr/main.f90 b/test/smoke-fort-dev/use-omp-get-mapped-ptr/main.f90 new file mode 100644 index 000000000..673981cbf --- /dev/null +++ b/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 + diff --git a/test/smoke-fort-fails/use-device-ptr-cptr/Makefile b/test/smoke-fort-fails/use-device-ptr-cptr/Makefile new file mode 100644 index 000000000..fb08569ec --- /dev/null +++ b/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 $@ diff --git a/test/smoke-fort-fails/use-device-ptr-cptr/bar.hip b/test/smoke-fort-fails/use-device-ptr-cptr/bar.hip new file mode 100644 index 000000000..659aa8828 --- /dev/null +++ b/test/smoke-fort-fails/use-device-ptr-cptr/bar.hip @@ -0,0 +1,23 @@ +#include +#include + +__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" */ diff --git a/test/smoke-fort-fails/use-device-ptr-cptr/main.f90 b/test/smoke-fort-fails/use-device-ptr-cptr/main.f90 new file mode 100644 index 000000000..0627a543d --- /dev/null +++ b/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 +