diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt index 31adbe45388a87..4e552e05895936 100644 --- a/offload/CMakeLists.txt +++ b/offload/CMakeLists.txt @@ -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() diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt index ae7d54d127edd5..0c9b77ded5abe0 100644 --- a/offload/DeviceRTL/CMakeLists.txt +++ b/offload/DeviceRTL/CMakeLists.txt @@ -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 ) diff --git a/offload/plugins-nextgen/CMakeLists.txt b/offload/plugins-nextgen/CMakeLists.txt index 5d684420c2901a..d31bf557669eac 100644 --- a/offload/plugins-nextgen/CMakeLists.txt +++ b/offload/plugins-nextgen/CMakeLists.txt @@ -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() diff --git a/offload/plugins-nextgen/amdgpu/CMakeLists.txt b/offload/plugins-nextgen/amdgpu/CMakeLists.txt index 4a79fa83420c8c..5f72568e9a537d 100644 --- a/offload/plugins-nextgen/amdgpu/CMakeLists.txt +++ b/offload/plugins-nextgen/amdgpu/CMakeLists.txt @@ -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) diff --git a/offload/plugins-nextgen/common/CMakeLists.txt b/offload/plugins-nextgen/common/CMakeLists.txt index 43e08d575a4ac7..284f98875170cd 100644 --- a/offload/plugins-nextgen/common/CMakeLists.txt +++ b/offload/plugins-nextgen/common/CMakeLists.txt @@ -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} diff --git a/offload/plugins-nextgen/cuda/CMakeLists.txt b/offload/plugins-nextgen/cuda/CMakeLists.txt index 4be33c6b69dced..51507e7537a422 100644 --- a/offload/plugins-nextgen/cuda/CMakeLists.txt +++ b/offload/plugins-nextgen/cuda/CMakeLists.txt @@ -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) diff --git a/offload/plugins-nextgen/host/CMakeLists.txt b/offload/plugins-nextgen/host/CMakeLists.txt index 72b5681283fe28..d35aadc56aa80d 100644 --- a/offload/plugins-nextgen/host/CMakeLists.txt +++ b/offload/plugins-nextgen/host/CMakeLists.txt @@ -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() diff --git a/offload/src/CMakeLists.txt b/offload/src/CMakeLists.txt index 590fb650c0ae88..c54bcbe7ed26fe 100644 --- a/offload/src/CMakeLists.txt +++ b/offload/src/CMakeLists.txt @@ -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} ) diff --git a/offload/test/mapping/array_section_implicit_capture.c b/offload/test/mapping/array_section_implicit_capture.c index 210b7e51cbdff3..1042bd23bd5efa 100644 --- a/offload/test/mapping/array_section_implicit_capture.c +++ b/offload/test/mapping/array_section_implicit_capture.c @@ -1,4 +1,4 @@ -// RUN: %libomptarget-compile-generic +// RUN: %libomptarget-compile-generic // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic diff --git a/offload/test/offloading/default_thread_limit.c b/offload/test/offloading/default_thread_limit.c index 4da02bbb152e60..d9be843dfbda49 100644 --- a/offload/test/offloading/default_thread_limit.c +++ b/offload/test/offloading/default_thread_limit.c @@ -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(); diff --git a/offload/test/offloading/fortran/basic-target-region-1D-array-section.f90 b/offload/test/offloading/fortran/basic-target-region-1D-array-section.f90 index 476b77e4a549bb..e0b0f2ea734e94 100644 --- a/offload/test/offloading/fortran/basic-target-region-1D-array-section.f90 +++ b/offload/test/offloading/fortran/basic-target-region-1D-array-section.f90 @@ -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 diff --git a/offload/test/offloading/fortran/basic-target-region-3D-array-section.f90 b/offload/test/offloading/fortran/basic-target-region-3D-array-section.f90 index 229798b57477dd..1b3f86c0dbfaa4 100644 --- a/offload/test/offloading/fortran/basic-target-region-3D-array-section.f90 +++ b/offload/test/offloading/fortran/basic-target-region-3D-array-section.f90 @@ -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 diff --git a/offload/test/offloading/fortran/double-target-call-with-declare-target.f90 b/offload/test/offloading/fortran/double-target-call-with-declare-target.f90 index 56c96727d47522..38cb055bfa2da6 100644 --- a/offload/test/offloading/fortran/double-target-call-with-declare-target.f90 +++ b/offload/test/offloading/fortran/double-target-call-with-declare-target.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-allocatable-map-scopes.f90 b/offload/test/offloading/fortran/target-map-allocatable-map-scopes.f90 index bb47d3de96d2a9..506976b339dfc7 100644 --- a/offload/test/offloading/fortran/target-map-allocatable-map-scopes.f90 +++ b/offload/test/offloading/fortran/target-map-allocatable-map-scopes.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-derived-type-full-1.f90 b/offload/test/offloading/fortran/target-map-derived-type-full-1.f90 index cb03708554fed0..f4e11f8b9e326c 100644 --- a/offload/test/offloading/fortran/target-map-derived-type-full-1.f90 +++ b/offload/test/offloading/fortran/target-map-derived-type-full-1.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-derived-type-full-2.f90 b/offload/test/offloading/fortran/target-map-derived-type-full-2.f90 index 0095b0fdf86a66..f97880a2023d4a 100644 --- a/offload/test/offloading/fortran/target-map-derived-type-full-2.f90 +++ b/offload/test/offloading/fortran/target-map-derived-type-full-2.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-derived-type-full-implicit-1.f90 b/offload/test/offloading/fortran/target-map-derived-type-full-implicit-1.f90 index f57e2c70d155e9..0f1d58de961cac 100644 --- a/offload/test/offloading/fortran/target-map-derived-type-full-implicit-1.f90 +++ b/offload/test/offloading/fortran/target-map-derived-type-full-implicit-1.f90 @@ -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 @@ end program main !CHECK: 10 !CHECK: 2. !CHECK: (2.,10.) - \ No newline at end of file diff --git a/offload/test/offloading/fortran/target-map-derived-type-full-implicit-2.f90 b/offload/test/offloading/fortran/target-map-derived-type-full-implicit-2.f90 index 92d3454d462a7d..a20b644652f823 100644 --- a/offload/test/offloading/fortran/target-map-derived-type-full-implicit-2.f90 +++ b/offload/test/offloading/fortran/target-map-derived-type-full-implicit-2.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-double-large-nested-dtype-multi-member.f90 b/offload/test/offloading/fortran/target-map-double-large-nested-dtype-multi-member.f90 index 31774be1914638..17e1cd9549f074 100644 --- a/offload/test/offloading/fortran/target-map-double-large-nested-dtype-multi-member.f90 +++ b/offload/test/offloading/fortran/target-map-double-large-nested-dtype-multi-member.f90 @@ -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. diff --git a/offload/test/offloading/fortran/target-map-double-nested-dtype-array-bounds.f90 b/offload/test/offloading/fortran/target-map-double-nested-dtype-array-bounds.f90 index cecfb9e84a59d6..d327ebf8c1a921 100644 --- a/offload/test/offloading/fortran/target-map-double-nested-dtype-array-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-double-nested-dtype-array-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-double-nested-dtype-double-array-bounds.f90 b/offload/test/offloading/fortran/target-map-double-nested-dtype-double-array-bounds.f90 index a8762a0829cc1d..d80ec8a4e0699d 100644 --- a/offload/test/offloading/fortran/target-map-double-nested-dtype-double-array-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-double-nested-dtype-double-array-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-double-nested-dtype-single-member.f90 b/offload/test/offloading/fortran/target-map-double-nested-dtype-single-member.f90 index 9ecb394dbe4629..59bcd4cc383405 100644 --- a/offload/test/offloading/fortran/target-map-double-nested-dtype-single-member.f90 +++ b/offload/test/offloading/fortran/target-map-double-nested-dtype-single-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit-update.f90 b/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit-update.f90 index 3b3ec96b9babfd..47fb6d210e9786 100644 --- a/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit-update.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit-update.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit.f90 b/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit.f90 index 5f7e9f94682659..e92e85ed0bc003 100644 --- a/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-arr-bounds-member-enter-exit.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-explicit-individual-array-member.f90 b/offload/test/offloading/fortran/target-map-dtype-explicit-individual-array-member.f90 index 907b16ffedf524..97b804e6d1d9d6 100644 --- a/offload/test/offloading/fortran/target-map-dtype-explicit-individual-array-member.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-explicit-individual-array-member.f90 @@ -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 @@ program main 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 diff --git a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-3D-member-bounds.f90 b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-3D-member-bounds.f90 index 110fb648980cdc..44727bec73c0e6 100644 --- a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-3D-member-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-3D-member-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member-bounds.f90 b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member-bounds.f90 index b7f6e2ddfb3bcb..fddc956e3fd534 100644 --- a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member.f90 b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member.f90 index c44a58dbebc8e4..2f74b28a67f6b7 100644 --- a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-array-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-member.f90 b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-member.f90 index a4205bc0179e20..9f2c185aa1ad09 100644 --- a/offload/test/offloading/fortran/target-map-dtype-multi-explicit-member.f90 +++ b/offload/test/offloading/fortran/target-map-dtype-multi-explicit-member.f90 @@ -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. diff --git a/offload/test/offloading/fortran/target-map-enter-exit-allocatables.f90 b/offload/test/offloading/fortran/target-map-enter-exit-allocatables.f90 index 865be95ba9682d..c14b23f6170d51 100644 --- a/offload/test/offloading/fortran/target-map-enter-exit-allocatables.f90 +++ b/offload/test/offloading/fortran/target-map-enter-exit-allocatables.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-enter-exit-array-2.f90 b/offload/test/offloading/fortran/target-map-enter-exit-array-2.f90 index 8d35a281caf933..fe65d6a00cbb8d 100644 --- a/offload/test/offloading/fortran/target-map-enter-exit-array-2.f90 +++ b/offload/test/offloading/fortran/target-map-enter-exit-array-2.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-enter-exit-array-bounds.f90 b/offload/test/offloading/fortran/target-map-enter-exit-array-bounds.f90 index d842cd15c65a10..dfb5d59273c616 100644 --- a/offload/test/offloading/fortran/target-map-enter-exit-array-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-enter-exit-array-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-enter-exit-array.f90 b/offload/test/offloading/fortran/target-map-enter-exit-array.f90 index 4a9fb6ee177f65..341c36647c4eda 100644 --- a/offload/test/offloading/fortran/target-map-enter-exit-array.f90 +++ b/offload/test/offloading/fortran/target-map-enter-exit-array.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-enter-exit-scalar.f90 b/offload/test/offloading/fortran/target-map-enter-exit-scalar.f90 index 70ed3d747829d9..f2d7099c041c15 100644 --- a/offload/test/offloading/fortran/target-map-enter-exit-scalar.f90 +++ b/offload/test/offloading/fortran/target-map-enter-exit-scalar.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-individual-dtype-member-map.f90 b/offload/test/offloading/fortran/target-map-individual-dtype-member-map.f90 index 4cdf41db70d43b..9a65053fd1455c 100644 --- a/offload/test/offloading/fortran/target-map-individual-dtype-member-map.f90 +++ b/offload/test/offloading/fortran/target-map-individual-dtype-member-map.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-large-nested-dtype-multi-member.f90 b/offload/test/offloading/fortran/target-map-large-nested-dtype-multi-member.f90 index 2412381e62e7fd..578b885046041c 100644 --- a/offload/test/offloading/fortran/target-map-large-nested-dtype-multi-member.f90 +++ b/offload/test/offloading/fortran/target-map-large-nested-dtype-multi-member.f90 @@ -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. diff --git a/offload/test/offloading/fortran/target-map-nested-dtype-complex-member.f90 b/offload/test/offloading/fortran/target-map-nested-dtype-complex-member.f90 index 9d29639ca101ee..e2e34e6b9b5c13 100644 --- a/offload/test/offloading/fortran/target-map-nested-dtype-complex-member.f90 +++ b/offload/test/offloading/fortran/target-map-nested-dtype-complex-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-nested-dtype-derived-member.f90 b/offload/test/offloading/fortran/target-map-nested-dtype-derived-member.f90 index 1be8cb32c8e6c8..b94915594a1562 100644 --- a/offload/test/offloading/fortran/target-map-nested-dtype-derived-member.f90 +++ b/offload/test/offloading/fortran/target-map-nested-dtype-derived-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-nested-dtype-multi-member.f90 b/offload/test/offloading/fortran/target-map-nested-dtype-multi-member.f90 index 6f4d5ad5c15bdb..c3d230853872d6 100644 --- a/offload/test/offloading/fortran/target-map-nested-dtype-multi-member.f90 +++ b/offload/test/offloading/fortran/target-map-nested-dtype-multi-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-nested-dtype-single-member.f90 b/offload/test/offloading/fortran/target-map-nested-dtype-single-member.f90 index 046fc13eb93c77..de436edb2ef24d 100644 --- a/offload/test/offloading/fortran/target-map-nested-dtype-single-member.f90 +++ b/offload/test/offloading/fortran/target-map-nested-dtype-single-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-pointer-scopes-enter-exit.f90 b/offload/test/offloading/fortran/target-map-pointer-scopes-enter-exit.f90 index dee75af06927be..e148fb123ecba6 100644 --- a/offload/test/offloading/fortran/target-map-pointer-scopes-enter-exit.f90 +++ b/offload/test/offloading/fortran/target-map-pointer-scopes-enter-exit.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-pointer-target-array-section-3d-bounds.f90 b/offload/test/offloading/fortran/target-map-pointer-target-array-section-3d-bounds.f90 index ff2298cf5dbc9d..c20fbbc0ff35cc 100644 --- a/offload/test/offloading/fortran/target-map-pointer-target-array-section-3d-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-pointer-target-array-section-3d-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-pointer-target-scopes.f90 b/offload/test/offloading/fortran/target-map-pointer-target-scopes.f90 index d9a7000719f0e8..4d5b995a7dfb3c 100644 --- a/offload/test/offloading/fortran/target-map-pointer-target-scopes.f90 +++ b/offload/test/offloading/fortran/target-map-pointer-target-scopes.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-dtype-explicit-member.f90 b/offload/test/offloading/fortran/target-map-two-dtype-explicit-member.f90 index b080b437e381b5..2ca621faf6ae2f 100644 --- a/offload/test/offloading/fortran/target-map-two-dtype-explicit-member.f90 +++ b/offload/test/offloading/fortran/target-map-two-dtype-explicit-member.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-dtype-individual-member-array-1D-bounds.f90 b/offload/test/offloading/fortran/target-map-two-dtype-individual-member-array-1D-bounds.f90 index 7cb7846ed0b241..77a5e1bc9e5fd1 100644 --- a/offload/test/offloading/fortran/target-map-two-dtype-individual-member-array-1D-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-two-dtype-individual-member-array-1D-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-1.f90 b/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-1.f90 index fbe6b305464a64..8675f470e1cb2c 100644 --- a/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-1.f90 +++ b/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-1.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-2.f90 b/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-2.f90 index 503329d90628fc..d13d9e8068ec22 100644 --- a/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-2.f90 +++ b/offload/test/offloading/fortran/target-map-two-dtype-mixed-implicit-explicit-capture-2.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-dtype-multi-member-array-1D-bounds.f90 b/offload/test/offloading/fortran/target-map-two-dtype-multi-member-array-1D-bounds.f90 index ed350c54dcb259..b788c6c5bcf262 100644 --- a/offload/test/offloading/fortran/target-map-two-dtype-multi-member-array-1D-bounds.f90 +++ b/offload/test/offloading/fortran/target-map-two-dtype-multi-member-array-1D-bounds.f90 @@ -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 diff --git a/offload/test/offloading/fortran/target-map-two-nested-dtype-member-array-map.f90 b/offload/test/offloading/fortran/target-map-two-nested-dtype-member-array-map.f90 index 42d9197f6e1e7a..2114a33e7ba129 100644 --- a/offload/test/offloading/fortran/target-map-two-nested-dtype-member-array-map.f90 +++ b/offload/test/offloading/fortran/target-map-two-nested-dtype-member-array-map.f90 @@ -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 diff --git a/offload/test/offloading/spmdization.c b/offload/test/offloading/spmdization.c index 77913bec8342f9..ab26456823e4bb 100644 --- a/offload/test/offloading/spmdization.c +++ b/offload/test/offloading/spmdization.c @@ -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; } diff --git a/offload/test/offloading/thread_limit.c b/offload/test/offloading/thread_limit.c index a8cc51b651dc96..f18a306614c790 100644 --- a/offload/test/offloading/thread_limit.c +++ b/offload/test/offloading/thread_limit.c @@ -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++) {