mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-18 12:16:49 +00:00
[Offload][NFC] Remove all trailing whitespace from offload/ (#92578)
Summary: This patch cleans up the training whitespace in a bunch of tests and CMake files. Most just in preparation for other cleanups.
This commit is contained in:
Joseph Huber
GitHub
parent
487d5af648
commit
16bb7e89a9
@ -134,7 +134,7 @@ if (NOT LIBOMPTARGET_LLVM_INCLUDE_DIRS)
|
||||
message(FATAL_ERROR "Missing definition for LIBOMPTARGET_LLVM_INCLUDE_DIRS")
|
||||
endif()
|
||||
|
||||
if(DEFINED LIBOMPTARGET_BUILD_CUDA_PLUGIN OR
|
||||
if(DEFINED LIBOMPTARGET_BUILD_CUDA_PLUGIN OR
|
||||
DEFINED LIBOMPTARGET_BUILD_AMDGPU_PLUGIN)
|
||||
message(WARNING "Option removed, use 'LIBOMPTARGET_PLUGINS_TO_BUILD' instead")
|
||||
endif()
|
||||
|
||||
@ -104,7 +104,7 @@ set(src_files
|
||||
# vectorized accesses to the shared state. Generally, those are "good" but
|
||||
# the optimizer pipeline (esp. Attributor) does not fully support vectorized
|
||||
# instructions yet and we end up missing out on way more important constant
|
||||
# propagation. That said, we will run the vectorizer again after the runtime
|
||||
# propagation. That said, we will run the vectorizer again after the runtime
|
||||
# has been linked into the user program.
|
||||
set(clang_opt_flags -O3 -mllvm -openmp-opt-disable -DSHARED_SCRATCHPAD_SIZE=512 -mllvm -vectorize-slp=false )
|
||||
set(link_opt_flags -O3 -openmp-opt-disable -attributor-enable=module -vectorize-slp=false )
|
||||
|
||||
@ -38,7 +38,7 @@ function(add_target_library target_name lib_name)
|
||||
PluginCommon ${OPENMP_PTHREAD_LIB})
|
||||
|
||||
target_compile_definitions(${target_name} PRIVATE TARGET_NAME=${lib_name})
|
||||
target_compile_definitions(${target_name} PRIVATE
|
||||
target_compile_definitions(${target_name} PRIVATE
|
||||
DEBUG_PREFIX="TARGET ${lib_name} RTL")
|
||||
set_target_properties(${target_name} PROPERTIES POSITION_INDEPENDENT_CODE ON)
|
||||
endfunction()
|
||||
|
||||
@ -28,7 +28,7 @@ endif()
|
||||
option(LIBOMPTARGET_FORCE_AMDGPU_TESTS "Build AMDGPU libomptarget tests" OFF)
|
||||
if (LIBOMPTARGET_FOUND_AMDGPU_GPU OR LIBOMPTARGET_FORCE_AMDGPU_TESTS)
|
||||
# Report to the parent scope that we are building a plugin for amdgpu
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"${LIBOMPTARGET_SYSTEM_TARGETS} amdgcn-amd-amdhsa" PARENT_SCOPE)
|
||||
list(APPEND LIBOMPTARGET_TESTED_PLUGINS "omptarget.rtl.amdgpu")
|
||||
set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
|
||||
|
||||
@ -51,7 +51,7 @@ target_compile_definitions(PluginCommon PRIVATE
|
||||
target_compile_options(PluginCommon PUBLIC ${offload_compile_flags})
|
||||
target_link_options(PluginCommon PUBLIC ${offload_link_flags})
|
||||
|
||||
target_include_directories(PluginCommon PUBLIC
|
||||
target_include_directories(PluginCommon PUBLIC
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/include
|
||||
${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
|
||||
${LIBOMPTARGET_BINARY_INCLUDE_DIR}
|
||||
|
||||
@ -25,7 +25,7 @@ endif()
|
||||
option(LIBOMPTARGET_FORCE_NVIDIA_TESTS "Build NVIDIA libomptarget tests" OFF)
|
||||
if (LIBOMPTARGET_FOUND_NVIDIA_GPU OR LIBOMPTARGET_FORCE_NVIDIA_TESTS)
|
||||
libomptarget_say("Enable tests using CUDA plugin")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"${LIBOMPTARGET_SYSTEM_TARGETS} nvptx64-nvidia-cuda nvptx64-nvidia-cuda-LTO" PARENT_SCOPE)
|
||||
list(APPEND LIBOMPTARGET_TESTED_PLUGINS "omptarget.rtl.cuda")
|
||||
set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)
|
||||
|
||||
@ -44,23 +44,23 @@ endif()
|
||||
|
||||
# Define the target specific triples and ELF machine values.
|
||||
if(CMAKE_SYSTEM_PROCESSOR MATCHES "ppc64le$")
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"powerpc64le-ibm-linux-gnu" "powerpc64le-ibm-linux-gnu-LTO")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
|
||||
elseif(CMAKE_SYSTEM_PROCESSOR MATCHES "ppc64$")
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"powerpc64-ibm-linux-gnu" "powerpc64-ibm-linux-gnu-LTO")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
|
||||
elseif(CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64$")
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"x86_64-pc-linux-gnu" "x86_64-pc-linux-gnu-LTO")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
|
||||
elseif(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64$")
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"aarch64-unknown-linux-gnu" "aarch64-unknown-linux-gnu-LTO")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
|
||||
elseif(CMAKE_SYSTEM_PROCESSOR MATCHES "s390x$")
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
list(APPEND LIBOMPTARGET_SYSTEM_TARGETS
|
||||
"s390x-ibm-linux-gnu" "s390x-ibm-linux-gnu-LTO")
|
||||
set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
|
||||
endif()
|
||||
|
||||
@ -38,7 +38,7 @@ add_llvm_library(omptarget
|
||||
NO_INSTALL_RPATH
|
||||
BUILDTREE_ONLY
|
||||
)
|
||||
target_include_directories(omptarget PRIVATE
|
||||
target_include_directories(omptarget PRIVATE
|
||||
${LIBOMPTARGET_INCLUDE_DIR} ${LIBOMPTARGET_BINARY_INCLUDE_DIR}
|
||||
)
|
||||
|
||||
|
||||
@ -1,4 +1,4 @@
|
||||
// RUN: %libomptarget-compile-generic
|
||||
// RUN: %libomptarget-compile-generic
|
||||
// RUN: %libomptarget-run-generic 2>&1 \
|
||||
// RUN: | %fcheck-generic
|
||||
|
||||
|
||||
@ -33,7 +33,7 @@ int main() {
|
||||
optnone();
|
||||
}
|
||||
// DEFAULT: [[NT:(128|256)]] (MaxFlatWorkGroupSize: [[NT]]
|
||||
#pragma omp target
|
||||
#pragma omp target
|
||||
#pragma omp teams distribute parallel for
|
||||
for (int i = 0; i < N; ++i) {
|
||||
optnone();
|
||||
|
||||
@ -1,4 +1,4 @@
|
||||
! Basic offloading test of arrays with provided lower
|
||||
! Basic offloading test of arrays with provided lower
|
||||
! and upper bounds as specified by OpenMP's sectioning
|
||||
! REQUIRES: flang
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -20,7 +20,7 @@ program main
|
||||
i = i + 1
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, write_arr(:)
|
||||
end program
|
||||
|
||||
|
||||
@ -12,7 +12,7 @@ program main
|
||||
implicit none
|
||||
integer :: inArray(3,3,3)
|
||||
integer :: outArray(3,3,3)
|
||||
integer :: i, j, k
|
||||
integer :: i, j, k
|
||||
integer :: j2 = 3, k2 = 3
|
||||
|
||||
do i = 1, 3
|
||||
@ -25,7 +25,7 @@ program main
|
||||
end do
|
||||
|
||||
j = 1
|
||||
k = 1
|
||||
k = 1
|
||||
!$omp target map(tofrom:inArray(1:3, 1:3, 2:2), outArray(1:3, 1:3, 1:3), j, k, j2, k2)
|
||||
do while (j <= j2)
|
||||
k = 1
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test with two target regions mapping the same
|
||||
! declare target Fortran array and writing some values to
|
||||
! it before checking the host correctly receives the
|
||||
! declare target Fortran array and writing some values to
|
||||
! it before checking the host correctly receives the
|
||||
! correct updates made on the device.
|
||||
! REQUIRES: flang
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -35,7 +35,7 @@ program main
|
||||
i = i + 1
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, sp(:)
|
||||
|
||||
end program
|
||||
|
||||
@ -11,55 +11,55 @@
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
module test
|
||||
contains
|
||||
subroutine func_arg(arg_alloc)
|
||||
subroutine func_arg(arg_alloc)
|
||||
integer, allocatable, intent (inout) :: arg_alloc(:)
|
||||
|
||||
!$omp target map(tofrom: arg_alloc)
|
||||
|
||||
!$omp target map(tofrom: arg_alloc)
|
||||
do index = 1, 10
|
||||
arg_alloc(index) = arg_alloc(index) + index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, arg_alloc
|
||||
end subroutine func_arg
|
||||
end module
|
||||
|
||||
subroutine func
|
||||
integer, allocatable :: local_alloc(:)
|
||||
end module
|
||||
|
||||
subroutine func
|
||||
integer, allocatable :: local_alloc(:)
|
||||
allocate(local_alloc(10))
|
||||
|
||||
!$omp target map(tofrom: local_alloc)
|
||||
|
||||
!$omp target map(tofrom: local_alloc)
|
||||
do index = 1, 10
|
||||
local_alloc(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, local_alloc
|
||||
|
||||
|
||||
deallocate(local_alloc)
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, allocatable :: map_ptr(:)
|
||||
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, allocatable :: map_ptr(:)
|
||||
|
||||
allocate(map_ptr(10))
|
||||
|
||||
!$omp target map(tofrom: map_ptr)
|
||||
|
||||
!$omp target map(tofrom: map_ptr)
|
||||
do index = 1, 10
|
||||
map_ptr(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
call func
|
||||
call func
|
||||
|
||||
print *, map_ptr
|
||||
|
||||
call func_arg(map_ptr)
|
||||
|
||||
deallocate(map_ptr)
|
||||
end program
|
||||
end program
|
||||
|
||||
! CHECK: 1 2 3 4 5 6 7 8 9 10
|
||||
! CHECK: 1 2 3 4 5 6 7 8 9 10
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type mapping when mapped
|
||||
! explicit derived type mapping when mapped
|
||||
! to target and assinging one derived type
|
||||
! to another
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
@ -16,22 +16,22 @@ program main
|
||||
integer(4) :: ix = 0
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: in
|
||||
type(scalar) :: out
|
||||
in%ix = 10
|
||||
in%rx = 2.0
|
||||
in%zx = (2, 10)
|
||||
|
||||
|
||||
!$omp target map(from:out) map(to:in)
|
||||
out = in
|
||||
out = in
|
||||
!$omp end target
|
||||
|
||||
|
||||
print*, in%ix
|
||||
print*, in%rx
|
||||
write (*,*) in%zx
|
||||
|
||||
|
||||
print*, out%ix
|
||||
print*, out%rx
|
||||
write (*,*) out%zx
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type mapping when mapped to
|
||||
! explicit derived type mapping when mapped to
|
||||
! target and assigning to individual members
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -16,29 +16,29 @@ program main
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
integer(4) :: array(5)
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: out
|
||||
type(scalar) :: in
|
||||
|
||||
|
||||
in%ix = 10
|
||||
in%rx = 2.0
|
||||
in%zx = (2, 10)
|
||||
|
||||
|
||||
do i = 1, 5
|
||||
in%array(i) = i
|
||||
end do
|
||||
|
||||
end do
|
||||
|
||||
!$omp target map(from:out) map(to:in)
|
||||
out%ix = in%ix
|
||||
out%rx = in%rx
|
||||
out%zx = in%zx
|
||||
|
||||
|
||||
do i = 1, 5
|
||||
out%array(i) = in%array(i)
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print*, in%ix
|
||||
print*, in%rx
|
||||
print*, in%array
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! implicit derived type mapping when mapped
|
||||
! implicit derived type mapping when mapped
|
||||
! to target and assinging one derived type
|
||||
! to another
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
@ -16,18 +16,18 @@ program main
|
||||
integer(4) :: ix = 0
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: in
|
||||
type(scalar) :: out
|
||||
in%ix = 10
|
||||
in%rx = 2.0
|
||||
in%zx = (2, 10)
|
||||
|
||||
|
||||
!$omp target map(from:out)
|
||||
out = in
|
||||
out = in
|
||||
!$omp end target
|
||||
|
||||
|
||||
print*, in%ix
|
||||
print*, in%rx
|
||||
write (*,*) in%zx
|
||||
@ -43,4 +43,3 @@ program main
|
||||
!CHECK: 10
|
||||
!CHECK: 2.
|
||||
!CHECK: (2.,10.)
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type mapping when mapped
|
||||
! explicit derived type mapping when mapped
|
||||
! to target and assinging one derived type
|
||||
! to another
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
@ -17,29 +17,29 @@ program main
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
integer(4) :: array(5)
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: out
|
||||
type(scalar) :: in
|
||||
|
||||
|
||||
in%ix = 10
|
||||
in%rx = 2.0
|
||||
in%zx = (2, 10)
|
||||
|
||||
|
||||
do i = 1, 5
|
||||
in%array(i) = i
|
||||
end do
|
||||
|
||||
end do
|
||||
|
||||
!$omp target
|
||||
out%ix = in%ix
|
||||
out%rx = in%rx
|
||||
out%zx = in%zx
|
||||
|
||||
|
||||
do i = 1, 5
|
||||
out%array(i) = in%array(i)
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print*, in%ix
|
||||
print*, in%rx
|
||||
print*, in%array
|
||||
|
||||
@ -21,14 +21,14 @@ program main
|
||||
integer(4) :: i3
|
||||
integer(4) :: j3
|
||||
integer(4) :: k3
|
||||
end type bottom_layer2
|
||||
end type bottom_layer2
|
||||
|
||||
type :: middle_layer
|
||||
real(4) :: array_i2(10)
|
||||
real(4) :: i2
|
||||
real(4) :: array_j2(10)
|
||||
type(bottom_layer1) :: nest
|
||||
type(bottom_layer2) :: nest2
|
||||
type(bottom_layer1) :: nest
|
||||
type(bottom_layer2) :: nest2
|
||||
end type middle_layer
|
||||
|
||||
type :: top_layer
|
||||
@ -39,18 +39,18 @@ program main
|
||||
integer(4) :: k
|
||||
type(middle_layer) :: nested
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
type(top_layer) :: top_dtype2
|
||||
|
||||
top_dtype2%nested%nest%i4 = 10
|
||||
top_dtype2%nested%nest%j4 = 12
|
||||
top_dtype2%nested%nest%k4 = 54
|
||||
|
||||
|
||||
top_dtype2%nested%nest2%i3 = 20
|
||||
top_dtype2%nested%nest2%j3 = 40
|
||||
top_dtype2%nested%nest2%k3 = 60
|
||||
|
||||
|
||||
top_dtype2%nested%i2 = 200
|
||||
|
||||
do i = 1, 10
|
||||
@ -64,20 +64,20 @@ program main
|
||||
!$omp map(to: top_dtype2%array_i, top_dtype2%nested%nest2%i3, top_dtype2%nested%i2) &
|
||||
!$omp map(to: top_dtype2%nested%nest2%k3, top_dtype2%nested%nest2%j3)
|
||||
top_dtype%nested%nest%i4 = top_dtype2%nested%nest%i4
|
||||
top_dtype%nested%nest%j4 = top_dtype2%nested%nest%j4
|
||||
top_dtype%nested%nest%j4 = top_dtype2%nested%nest%j4
|
||||
top_dtype%nested%nest%k4 = top_dtype2%nested%nest%k4
|
||||
|
||||
|
||||
top_dtype%nested%nest2%i3 = top_dtype2%nested%nest2%i3
|
||||
top_dtype%nested%nest2%j3 = top_dtype2%nested%nest2%j3
|
||||
top_dtype%nested%nest2%k3 = top_dtype2%nested%nest2%k3
|
||||
|
||||
|
||||
top_dtype%nested%i2 = top_dtype2%nested%i2
|
||||
|
||||
do i = 1, 10
|
||||
top_dtype%array_i(i) = top_dtype2%array_i(i)
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%nest%i4
|
||||
print *, top_dtype%nested%nest%j4
|
||||
print *, top_dtype%nested%nest%k4
|
||||
@ -85,10 +85,10 @@ program main
|
||||
print *, top_dtype%nested%nest2%i3
|
||||
print *, top_dtype%nested%nest2%j3
|
||||
print *, top_dtype%nested%nest2%k3
|
||||
|
||||
|
||||
print *, top_dtype%nested%i2
|
||||
|
||||
print *, top_dtype%array_i
|
||||
print *, top_dtype%array_i
|
||||
end program main
|
||||
|
||||
!CHECK: 10.
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of two
|
||||
! explicit arrau member maps with bounds from
|
||||
! two nested derived types
|
||||
! explicit arrau member maps with bounds from
|
||||
! two nested derived types
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -25,20 +25,20 @@ program main
|
||||
integer, allocatable :: array_j(:)
|
||||
integer(4) :: k
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
type(top_layer) :: top_dtype2
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%array_i2(4:8), top_dtype2%nested%array_j2(4:8))
|
||||
do i = 4, 8
|
||||
do i = 4, 8
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
do i = 4, 8
|
||||
do i = 4, 8
|
||||
top_dtype2%nested%array_j2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%array_i2
|
||||
print *, top_dtype2%nested%array_j2
|
||||
end program main
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of two
|
||||
! explicit array member maps with array bounds
|
||||
! from two nested derived types
|
||||
! explicit array member maps with array bounds
|
||||
! from two nested derived types
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -25,20 +25,20 @@ program main
|
||||
integer, allocatable :: array_j(:)
|
||||
integer(4) :: k
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
type(top_layer) :: top_dtype2
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%array_i2(4:8), top_dtype2%nested%array_j2(4:8))
|
||||
do i = 4, 8
|
||||
do i = 4, 8
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
do i = 4, 8
|
||||
do i = 4, 8
|
||||
top_dtype2%nested%array_j2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%array_i2
|
||||
print *, top_dtype2%nested%array_j2
|
||||
end program main
|
||||
|
||||
@ -25,18 +25,18 @@ program main
|
||||
integer, allocatable :: array_j(:)
|
||||
integer(4) :: k
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
type(top_layer) :: top_dtype2
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%array_i2, top_dtype2%nested%array_j2)
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype2%nested%array_j2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
print *, top_dtype%nested%array_i2
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type member mapping of
|
||||
! an array with bounds when mapped to
|
||||
! explicit derived type member mapping of
|
||||
! an array with bounds when mapped to
|
||||
! target using a combination of update,
|
||||
! enter and exit directives.
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
@ -25,24 +25,24 @@ program main
|
||||
|
||||
!$omp target enter data map(to: scalar_arr%array(3:6))
|
||||
|
||||
! overwrite our target data with an update.
|
||||
! overwrite our target data with an update.
|
||||
do I = 1, 10
|
||||
scalar_arr%array(I) = 10
|
||||
end do
|
||||
|
||||
!$omp target update to(scalar_arr%array(3:6))
|
||||
|
||||
! The compiler/runtime is less friendly about read/write out of
|
||||
! The compiler/runtime is less friendly about read/write out of
|
||||
! bounds when using enter and exit, we have to specifically loop
|
||||
! over the correct range
|
||||
!$omp target
|
||||
do i=3,6
|
||||
scalar_arr%array(i) = scalar_arr%array(i) + i
|
||||
end do
|
||||
!$omp end target
|
||||
!$omp end target
|
||||
|
||||
!$omp target exit data map(from: scalar_arr%array(3:6))
|
||||
|
||||
|
||||
print*, scalar_arr%array
|
||||
end program
|
||||
|
||||
|
||||
@ -1,7 +1,7 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type member mapping of
|
||||
! an array with bounds when mapped to
|
||||
! target using a combination of enter and
|
||||
! explicit derived type member mapping of
|
||||
! an array with bounds when mapped to
|
||||
! target using a combination of enter and
|
||||
! exit directives.
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -24,25 +24,25 @@ program main
|
||||
end do
|
||||
|
||||
!$omp target enter data map(to: scalar_arr%array(3:6))
|
||||
|
||||
|
||||
! Shouldn't overwrite data already locked in
|
||||
! on target via enter, which will then be
|
||||
! on target via enter, which will then be
|
||||
! overwritten by our exit
|
||||
do I = 1, 10
|
||||
scalar_arr%array(I) = 10
|
||||
end do
|
||||
|
||||
! The compiler/runtime is less friendly about read/write out of
|
||||
! The compiler/runtime is less friendly about read/write out of
|
||||
! bounds when using enter and exit, we have to specifically loop
|
||||
! over the correct range
|
||||
!$omp target
|
||||
do i=3,6
|
||||
scalar_arr%array(i) = scalar_arr%array(i) + i
|
||||
end do
|
||||
!$omp end target
|
||||
!$omp end target
|
||||
|
||||
!$omp target exit data map(from: scalar_arr%array(3:6))
|
||||
|
||||
|
||||
print*, scalar_arr%array
|
||||
end program
|
||||
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type member mapping of
|
||||
! explicit derived type member mapping of
|
||||
! an array when mapped to target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -18,9 +18,9 @@ type :: scalar_array
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr
|
||||
|
||||
|
||||
!$omp target map(tofrom:scalar_arr%array_y)
|
||||
do i = 1, 10
|
||||
scalar_arr%array_y(i) = i
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type member mapping of
|
||||
! two arrays with explicit bounds when
|
||||
! explicit derived type member mapping of
|
||||
! two arrays with explicit bounds when
|
||||
! mapped to target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -19,9 +19,9 @@ program main
|
||||
integer(4) :: array_y(3,3,3)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr
|
||||
|
||||
|
||||
do i = 1, 3
|
||||
do j = 1, 3
|
||||
do k = 1, 3
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit derived type member mapping of
|
||||
! two arrays with explicit bounds when
|
||||
! explicit derived type member mapping of
|
||||
! two arrays with explicit bounds when
|
||||
! mapped to target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -19,9 +19,9 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr%array_x(i) = i
|
||||
end do
|
||||
|
||||
@ -18,9 +18,9 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr%array_x(i) = i
|
||||
end do
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! derived type mapping of two explicit
|
||||
! derived type mapping of two explicit
|
||||
! members to target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -11,23 +11,23 @@
|
||||
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
program main
|
||||
type :: scalar
|
||||
integer(4) :: ix = 0
|
||||
type :: scalar
|
||||
integer(4) :: ix = 0
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
real(4) :: ry = 1.0
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: scalar_struct
|
||||
|
||||
|
||||
!$omp target map(from:scalar_struct%rx, scalar_struct%ry)
|
||||
scalar_struct%rx = 21.0
|
||||
scalar_struct%ry = 27.0
|
||||
!$omp end target
|
||||
|
||||
|
||||
print*, scalar_struct%rx
|
||||
print*, scalar_struct%ry
|
||||
end program main
|
||||
|
||||
|
||||
!CHECK: 21.
|
||||
!CHECK: 27.
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of allocatables
|
||||
! with enter, exit and target
|
||||
! with enter, exit and target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -10,9 +10,9 @@
|
||||
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
program main
|
||||
integer, allocatable :: A(:)
|
||||
integer, allocatable :: A(:)
|
||||
allocate(A(10))
|
||||
|
||||
|
||||
!$omp target enter data map(alloc: A)
|
||||
|
||||
!$omp target
|
||||
@ -24,11 +24,11 @@ program main
|
||||
!$omp target exit data map(from: A)
|
||||
|
||||
!$omp target exit data map(delete: A)
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
print *, A(i)
|
||||
end do
|
||||
|
||||
|
||||
deallocate(A)
|
||||
end program
|
||||
|
||||
|
||||
@ -18,7 +18,7 @@ program main
|
||||
|
||||
!$omp target enter data map(to: array)
|
||||
! Shouldn't overwrite data already locked in
|
||||
! on target via enter, this will then be
|
||||
! on target via enter, this will then be
|
||||
! overwritten by our exit
|
||||
do I = 1, 10
|
||||
array(I) = 10
|
||||
@ -28,7 +28,7 @@ program main
|
||||
do i=1,10
|
||||
array(i) = array(i) + i
|
||||
end do
|
||||
!$omp end target
|
||||
!$omp end target
|
||||
|
||||
!$omp target exit data map(from: array)
|
||||
print*, array
|
||||
|
||||
@ -20,20 +20,20 @@ program main
|
||||
|
||||
!$omp target enter data map(to: array(3:6))
|
||||
! Shouldn't overwrite data already locked in
|
||||
! on target via enter, which will then be
|
||||
! on target via enter, which will then be
|
||||
! overwritten by our exit
|
||||
do I = 1, 10
|
||||
array(I) = 10
|
||||
end do
|
||||
|
||||
! The compiler/runtime is less lenient about read/write out of
|
||||
! The compiler/runtime is less lenient about read/write out of
|
||||
! bounds when using enter and exit, we have to specifically loop
|
||||
! over the correctly mapped range
|
||||
!$omp target
|
||||
do i=3,6
|
||||
array(i) = array(i) + i
|
||||
end do
|
||||
!$omp end target
|
||||
!$omp end target
|
||||
|
||||
!$omp target exit data map(from: array(3:6))
|
||||
print *, array
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of fixed size
|
||||
! arrays with enter, exit and target
|
||||
! arrays with enter, exit and target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -11,7 +11,7 @@
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
program main
|
||||
integer :: A(10)
|
||||
|
||||
|
||||
!$omp target enter data map(alloc: A)
|
||||
|
||||
!$omp target
|
||||
@ -23,7 +23,7 @@ program main
|
||||
!$omp target exit data map(from: A)
|
||||
|
||||
!$omp target exit data map(delete: A)
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
print *, A(i)
|
||||
end do
|
||||
|
||||
@ -19,12 +19,12 @@ program main
|
||||
|
||||
!$omp target
|
||||
scalar = scalar + 50
|
||||
!$omp end target
|
||||
!$omp end target
|
||||
|
||||
!$omp target exit data map(from: scalar)
|
||||
|
||||
! not the answer one may expect, but it is the same
|
||||
! answer Clang gives so we are correctly on par with
|
||||
! not the answer one may expect, but it is the same
|
||||
! answer Clang gives so we are correctly on par with
|
||||
! Clang for the moment.
|
||||
print *, scalar
|
||||
end program
|
||||
|
||||
@ -17,8 +17,8 @@ program main
|
||||
real(4) :: rx = 0.0
|
||||
complex(4) :: zx = (0,0)
|
||||
real(4) :: ry = 1.0
|
||||
end type scalar
|
||||
|
||||
end type scalar
|
||||
|
||||
type(scalar) :: scalar_struct
|
||||
scalar_struct%rx = 2.0
|
||||
test = 21.0
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of an
|
||||
! explicit member map a large nested derived
|
||||
! explicit member map a large nested derived
|
||||
! type
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -21,14 +21,14 @@ program main
|
||||
integer(4) :: i3
|
||||
integer(4) :: j3
|
||||
integer(4) :: k3
|
||||
end type bottom_layer2
|
||||
end type bottom_layer2
|
||||
|
||||
type :: middle_layer
|
||||
real(4) :: array_i2(10)
|
||||
real(4) :: i2
|
||||
real(4) :: array_j2(10)
|
||||
type(bottom_layer1) :: nest
|
||||
type(bottom_layer2) :: nest2
|
||||
type(bottom_layer1) :: nest
|
||||
type(bottom_layer2) :: nest2
|
||||
end type middle_layer
|
||||
|
||||
type :: top_layer
|
||||
@ -39,7 +39,7 @@ program main
|
||||
integer(4) :: k
|
||||
type(middle_layer) :: nested
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
|
||||
top_dtype%nested%nest%j4 = 12
|
||||
@ -48,18 +48,18 @@ program main
|
||||
top_dtype%nested%nest%i4 = 10
|
||||
top_dtype%nested%nest%j4 = 12 + top_dtype%nested%nest%j4
|
||||
top_dtype%nested%nest%k4 = 54
|
||||
|
||||
|
||||
top_dtype%nested%nest2%i3 = 20
|
||||
top_dtype%nested%nest2%j3 = 40
|
||||
top_dtype%nested%nest2%k3 = 60
|
||||
|
||||
|
||||
top_dtype%nested%i2 = 200
|
||||
|
||||
do i = 1, 10
|
||||
top_dtype%array_i(i) = i
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%nest%i4
|
||||
print *, top_dtype%nested%nest%j4
|
||||
print *, top_dtype%nested%nest%k4
|
||||
@ -67,10 +67,10 @@ program main
|
||||
print *, top_dtype%nested%nest2%i3
|
||||
print *, top_dtype%nested%nest2%j3
|
||||
print *, top_dtype%nested%nest2%k3
|
||||
|
||||
|
||||
print *, top_dtype%nested%i2
|
||||
|
||||
print *, top_dtype%array_i
|
||||
print *, top_dtype%array_i
|
||||
end program main
|
||||
|
||||
!CHECK: 10.
|
||||
|
||||
@ -27,21 +27,21 @@ program main
|
||||
integer(4) :: k
|
||||
complex :: l
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%i2, top_dtype%k, top_dtype%nested%j2, top_dtype%nested%array_i2, top_dtype%l)
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
top_dtype%l = (10,20)
|
||||
top_dtype%nested%j2 = (510,210)
|
||||
|
||||
|
||||
top_dtype%nested%i2 = 30.30
|
||||
top_dtype%k = 74
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%i2
|
||||
print *, top_dtype%k
|
||||
print *, top_dtype%nested%array_i2
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of an
|
||||
! nested derived type member map with the
|
||||
! inclusion of an entire nested derived
|
||||
! nested derived type member map with the
|
||||
! inclusion of an entire nested derived
|
||||
! type being mapped
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -27,19 +27,19 @@ program main
|
||||
integer(4) :: k
|
||||
type(bottom_layer) :: nested2
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
|
||||
!$omp target map(tofrom: top_dtype%k, top_dtype%nested2%array_i2, top_dtype%nested)
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype%nested2%array_i2(i) = i * 2
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
top_dtype%nested%i2 = 30.30
|
||||
top_dtype%k = 74
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%i2
|
||||
print *, top_dtype%k
|
||||
print *, top_dtype%nested%array_i2
|
||||
|
||||
@ -25,18 +25,18 @@ program main
|
||||
integer, allocatable :: array_j(:)
|
||||
integer(4) :: k
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%i2, top_dtype%k, top_dtype%nested%array_i2)
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
|
||||
top_dtype%nested%i2 = 30.30
|
||||
top_dtype%k = 74
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%i2
|
||||
print *, top_dtype%k
|
||||
print *, top_dtype%nested%array_i2
|
||||
|
||||
@ -25,15 +25,15 @@ program main
|
||||
integer, allocatable :: array_j(:)
|
||||
integer(4) :: k
|
||||
end type top_layer
|
||||
|
||||
|
||||
type(top_layer) :: top_dtype
|
||||
|
||||
!$omp target map(tofrom: top_dtype%nested%array_i2)
|
||||
do i = 1, 10
|
||||
do i = 1, 10
|
||||
top_dtype%nested%array_i2(i) = i * 2
|
||||
end do
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, top_dtype%nested%array_i2
|
||||
end program main
|
||||
|
||||
|
||||
@ -11,52 +11,52 @@
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
module test
|
||||
contains
|
||||
subroutine func_arg(arg_alloc)
|
||||
subroutine func_arg(arg_alloc)
|
||||
integer, pointer, intent (inout) :: arg_alloc(:)
|
||||
|
||||
|
||||
!$omp target enter data map(alloc: arg_alloc)
|
||||
|
||||
!$omp target
|
||||
|
||||
!$omp target
|
||||
do index = 1, 10
|
||||
arg_alloc(index) = arg_alloc(index) + index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
!$omp target exit data map(from: arg_alloc)
|
||||
|
||||
|
||||
!$omp target exit data map(delete: arg_alloc)
|
||||
|
||||
|
||||
print *, arg_alloc
|
||||
end subroutine func_arg
|
||||
end module
|
||||
|
||||
subroutine func
|
||||
integer, pointer :: local_alloc(:)
|
||||
end module
|
||||
|
||||
subroutine func
|
||||
integer, pointer :: local_alloc(:)
|
||||
allocate(local_alloc(10))
|
||||
|
||||
!$omp target enter data map(alloc: local_alloc)
|
||||
|
||||
!$omp target
|
||||
!$omp target
|
||||
do index = 1, 10
|
||||
local_alloc(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
!$omp target exit data map(from: local_alloc)
|
||||
|
||||
!$omp target exit data map(delete: local_alloc)
|
||||
|
||||
print *, local_alloc
|
||||
|
||||
|
||||
deallocate(local_alloc)
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, pointer :: map_ptr(:)
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, pointer :: map_ptr(:)
|
||||
allocate(map_ptr(10))
|
||||
|
||||
|
||||
!$omp target enter data map(alloc: map_ptr)
|
||||
|
||||
!$omp target
|
||||
@ -64,12 +64,12 @@ program main
|
||||
map_ptr(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
!$omp target exit data map(from: map_ptr)
|
||||
|
||||
!$omp target exit data map(delete: map_ptr)
|
||||
|
||||
call func
|
||||
call func
|
||||
|
||||
print *, map_ptr
|
||||
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of pointer
|
||||
! and target with target where 3-D bounds have
|
||||
! and target with target where 3-D bounds have
|
||||
! been specified
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -15,10 +15,10 @@ program main
|
||||
integer, pointer :: outArray(:,:,:)
|
||||
integer, target :: in(3,3,3)
|
||||
integer, target :: out(3,3,3)
|
||||
|
||||
|
||||
inArray => in
|
||||
outArray => out
|
||||
|
||||
|
||||
do i = 1, 3
|
||||
do j = 1, 3
|
||||
do k = 1, 3
|
||||
|
||||
@ -1,5 +1,5 @@
|
||||
! Offloading test checking interaction of pointer
|
||||
! and target with target across multiple scopes
|
||||
! and target with target across multiple scopes
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
||||
@ -11,48 +11,48 @@
|
||||
! RUN: %libomptarget-compile-fortran-run-and-check-generic
|
||||
module test
|
||||
contains
|
||||
subroutine func_arg(arg_alloc)
|
||||
subroutine func_arg(arg_alloc)
|
||||
integer, pointer, intent (inout) :: arg_alloc(:)
|
||||
|
||||
!$omp target map(tofrom: arg_alloc)
|
||||
|
||||
!$omp target map(tofrom: arg_alloc)
|
||||
do index = 1, 10
|
||||
arg_alloc(index) = arg_alloc(index) + index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, arg_alloc
|
||||
end subroutine func_arg
|
||||
end module
|
||||
|
||||
subroutine func
|
||||
integer, pointer :: local_alloc(:)
|
||||
|
||||
subroutine func
|
||||
integer, pointer :: local_alloc(:)
|
||||
integer, target :: b(10)
|
||||
local_alloc => b
|
||||
|
||||
!$omp target map(tofrom: local_alloc)
|
||||
|
||||
!$omp target map(tofrom: local_alloc)
|
||||
do index = 1, 10
|
||||
local_alloc(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
|
||||
print *, local_alloc
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, pointer :: map_ptr(:)
|
||||
end subroutine func
|
||||
|
||||
|
||||
program main
|
||||
use test
|
||||
integer, pointer :: map_ptr(:)
|
||||
integer, target :: b(10)
|
||||
|
||||
|
||||
map_ptr => b
|
||||
|
||||
!$omp target map(tofrom: map_ptr)
|
||||
|
||||
!$omp target map(tofrom: map_ptr)
|
||||
do index = 1, 10
|
||||
map_ptr(index) = index
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
call func
|
||||
|
||||
call func
|
||||
|
||||
print *, map_ptr
|
||||
|
||||
|
||||
@ -18,10 +18,10 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
!$omp target map(tofrom:scalar_arr1%break_1, scalar_arr2%break_3)
|
||||
scalar_arr2%break_3 = 10
|
||||
scalar_arr1%break_1 = 15
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of two
|
||||
! derived type's with a single explicit array
|
||||
! member each being mapped with bounds to
|
||||
! member each being mapped with bounds to
|
||||
! target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -19,10 +19,10 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
|
||||
!$omp target map(tofrom:scalar_arr1%array_x(3:6), scalar_arr2%array_x(3:6))
|
||||
do i = 3, 6
|
||||
@ -31,7 +31,7 @@ program main
|
||||
end do
|
||||
!$omp end target
|
||||
|
||||
print*, scalar_arr1%array_x
|
||||
print*, scalar_arr1%array_x
|
||||
print*, scalar_arr2%array_x
|
||||
end program main
|
||||
|
||||
|
||||
@ -18,10 +18,10 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
!$omp target map(tofrom:scalar_arr1%break_1)
|
||||
scalar_arr2%break_3 = 10
|
||||
scalar_arr1%break_1 = 15
|
||||
|
||||
@ -19,13 +19,13 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr1%array_x(i) = i
|
||||
end do
|
||||
end do
|
||||
|
||||
!$omp target map(tofrom:scalar_arr2%array_x(3:6))
|
||||
do i = 3, 6
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
! Offloading test checking interaction of two
|
||||
! derived type's with two explicit array
|
||||
! members each being mapped with bounds to
|
||||
! members each being mapped with bounds to
|
||||
! target
|
||||
! REQUIRES: flang, amdgcn-amd-amdhsa
|
||||
! UNSUPPORTED: nvptx64-nvidia-cuda
|
||||
@ -19,10 +19,10 @@ program main
|
||||
real(4) :: array_y(10)
|
||||
real(4) :: break_3
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr1%array_x(i) = i
|
||||
scalar_arr2%array_x(i) = i
|
||||
@ -32,7 +32,7 @@ program main
|
||||
do i = 1, 10
|
||||
scalar_arr2%array_y(i) = scalar_arr1%array_x(i)
|
||||
end do
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr1%array_y(i) = scalar_arr2%array_x(i)
|
||||
end do
|
||||
|
||||
@ -24,10 +24,10 @@ program main
|
||||
real(4) :: break_3
|
||||
type(array) :: nested
|
||||
end type scalar_array
|
||||
|
||||
|
||||
type(scalar_array) :: scalar_arr1
|
||||
type(scalar_array) :: scalar_arr2
|
||||
|
||||
|
||||
do i = 1, 10
|
||||
scalar_arr1%nested%array_z(i) = i
|
||||
scalar_arr2%nested%array_z(i) = i
|
||||
@ -37,7 +37,7 @@ program main
|
||||
do i = 3, 6
|
||||
scalar_arr2%nested%array_ix(i) = scalar_arr1%nested%array_z(i)
|
||||
end do
|
||||
|
||||
|
||||
do i = 3, 6
|
||||
scalar_arr1%nested%array_ix(i) = scalar_arr2%nested%array_z(i)
|
||||
end do
|
||||
|
||||
@ -32,13 +32,19 @@ int main(void) {
|
||||
tid = omp_get_thread_num();
|
||||
maxt = omp_get_max_threads();
|
||||
#pragma omp parallel
|
||||
noop();
|
||||
noop();
|
||||
}
|
||||
printf("NumThreads: %i, InParallel: %i, Level: %i, ActiveLevel: %i, Nested: %i, "
|
||||
"ThreadNum: %i, MaxThreads: %i\n",
|
||||
nthreads, ip, lvl, alvl, nested, tid, maxt);
|
||||
// GENERIC: Generic mode
|
||||
// SPMD: Generic-SPMD mode
|
||||
// CHECK: NumThreads: 1, InParallel: 0, Level: 0, ActiveLevel: 0, Nested: 0, ThreadNum: 0, MaxThreads:
|
||||
// CHECK: NumThreads: 1
|
||||
// CHECK: InParallel: 0
|
||||
// CHECK: Level: 0
|
||||
// CHECK: ActiveLevel: 0
|
||||
// CHECK: Nested: 0
|
||||
// CHECK: ThreadNum: 0
|
||||
// CHECK: MaxThreads:
|
||||
return 0;
|
||||
}
|
||||
|
||||
@ -16,13 +16,13 @@ int main() {
|
||||
int n = 1 << 20;
|
||||
int th = 12;
|
||||
int te = n / th;
|
||||
// DEFAULT: 12 (MaxFlatWorkGroupSize:
|
||||
// DEFAULT: 12 (MaxFlatWorkGroupSize:
|
||||
#pragma omp target
|
||||
#pragma omp teams loop num_teams(te), thread_limit(th)
|
||||
for (int i = 0; i < n; i++) {
|
||||
}
|
||||
|
||||
// DEFAULT: 13 (MaxFlatWorkGroupSize:
|
||||
// DEFAULT: 13 (MaxFlatWorkGroupSize:
|
||||
#pragma omp target
|
||||
#pragma omp teams distribute parallel for simd num_teams(te), thread_limit(th+1) simdlen(64)
|
||||
for(int i = 0; i < n; i++) {
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user