From 6c6a2b184504a956a17600f4420fe7c3c6342f8e Mon Sep 17 00:00:00 2001 From: Michael Kruse Date: Thu, 18 Jul 2024 10:25:12 +0200 Subject: [PATCH 1/9] [sollve_vv] Autodetected native GPU (#129) Using the CMake option `-DTEST_SUITE_SYSTEM_GPU=native` or (`-DTEST_SUITE_SYSTEM_GPU=auto`), automatically pre-selects the tests that are expected to work based on the system's GPU. Also fixes a lot of other issues: * `SYSTEM_GPU` violates naming conventions * `SYSTEM_GPU` is not a cached variable, i.e. after re-running CMake (e.g. after changing a `CMakeLists.txt`), its value is forgotten * Remove `intel` and `regression` options. There are no lists for these * Remove possibility to specify multiple pre-select lists. The union of those would run, causing them to fail if run on the other vendor's GPU. * Remove duplicate `test_metadirective_arch_is_nvidia.c` in preselect lists * Update documentation * Add messages whenever SOLLVE V&V is skipped * Do not dump test directories into the builddir root * Replace `TEST_SUITE_FORCE_ALL` with `-DTEST_SUITE_SYSTEM_GPU=all` * Remove uses of `remove_definitions`. Just don't add those flags if you don't want them. The variable `SYSTEM_GPU` is also used in External/HeCBench and External/smoke. These are based on sollve_vv's previous `CMakeLists.txt` and are not affected. --- External/OpenMPOffloading.cmake | 78 +++++++++ External/sollve_vv/CMakeLists.txt | 149 ++++++------------ External/sollve_vv/README | 87 +++++----- External/sollve_vv/sollvevv_amd_tests.cmake | 10 +- .../sollve_vv/sollvevv_nvidia_tests.cmake | 11 +- 5 files changed, 178 insertions(+), 157 deletions(-) create mode 100644 External/OpenMPOffloading.cmake diff --git a/External/OpenMPOffloading.cmake b/External/OpenMPOffloading.cmake new file mode 100644 index 000000000..7a8e345df --- /dev/null +++ b/External/OpenMPOffloading.cmake @@ -0,0 +1,78 @@ +# Introspection of OpenMP offloading support +# TODO: Could become a find_package module + +if (NOT OpenMP_FOUND) + message(STATUS "OpenMP Offloading not available because OpenMP is not available") + return () +endif () + + +set(TEST_SUITE_SYSTEM_GPU "auto" CACHE STRING "Select only tests that Clang/Flang is expected to work with for the selected GPU (auto, native, amd, nvidia, or all)") +set_property(CACHE TEST_SUITE_SYSTEM_GPU PROPERTY STRINGS "amd" "nvidia" "all" "native" "auto") + +# TODO: Run only once if included multiple times +set(TEST_SUITE_EFFECTIVE_SYSTEM_GPU "${TEST_SUITE_SYSTEM_GPU}") +if (TEST_SUITE_EFFECTIVE_SYSTEM_GPU STREQUAL "auto") + # Test sets only maintained for Flang/Clang + if ((CMAKE_C_COMPILER_ID STREQUAL "Clang") OR (CMAKE_CXX_COMPILER_ID STREQUAL "Clang") OR (CMAKE_Fortran_COMPILER_ID STREQUAL "LLVMFlang")) + set (TEST_SUITE_EFFECTIVE_SYSTEM_GPU "native") + else () + set (TEST_SUITE_EFFECTIVE_SYSTEM_GPU "all") + endif () +endif () +if (TEST_SUITE_EFFECTIVE_SYSTEM_GPU STREQUAL "native") + # For LLVM, nvptx-arch and amdgpu-arch are executables in the bin directory, next to clang/flang. + set(_searchpaths) + foreach (_compiler IN ITEMS ${CMAKE_C_COMPILER} ${CMAKE_CXX_COMPILER} ${CMAKE_Fortran_COMPILER}) + get_filename_component(_compiler_dir "${_compiler}" DIRECTORY) + list(APPEND _searchpaths "${_compiler_dir}") + endforeach () + + find_program(NVPTX_ARCH_EXECUTABLE nvptx-arch HINTS ${_searchpaths}) + execute_process(COMMAND "${NVPTX_ARCH_EXECUTABLE}" RESULT_VARIABLE _nvptx_arch_result OUTPUT_VARIABLE _nvptx_arch_output ERROR_VARIABLE _nvptx_arch_error) + if (_nvptx_arch_result STREQUAL "0" AND _nvptx_arch_output) + # Which Nvidia GPU is detected currently does not matter + set(TEST_SUITE_EFFECTIVE_SYSTEM_GPU "nvidia") + endif () + + find_program(AMDGPU_ARCH_EXECUTABLE amdgpu-arch HINTS ${_searchpaths}) + execute_process(COMMAND "${AMDGPU_ARCH_EXECUTABLE}" RESULT_VARIABLE _amdgpu_arch_result OUTPUT_VARIABLE _amdgpu_arch_output ERROR_VARIABLE _amdgpu_arch_error) + if (_amdgpu_arch_result STREQUAL "0" AND _amdgpu_arch_output) + # Which AMD GPU is detected currently does not matter + set(TEST_SUITE_EFFECTIVE_SYSTEM_GPU "amd") + endif () + + if (TEST_SUITE_EFFECTIVE_SYSTEM_GPU STREQUAL "native") + if (TEST_SUITE_SYSTEM_GPU STREQUAL "auto") + # If no preselected list is available, run all tests + set(TEST_SUITE_EFFECTIVE_SYSTEM_GPU "all") + else () + message(STATUS "OpenMP Offloading not available because no native GPU detected") + return() + endif () + else () + message(STATUS "Native GPU is ${_effective_SYSTEM_GPU}") + endif () +endif () + + +# Do not try to create OpenMP_Offloading multiple times. +if (TARGET OpenMP_Offloading) + return () +endif () + +add_library(OpenMP_Offloading INTERFACE IMPORTED) +foreach (_lang IN ITEMS C CXX Fortran) + if (lang STREQUAL "Fortran" AND NOT TEST_SUITE_FORTRAN) + continue () + endif () + set(TEST_SUITE_OFFLOADING_${_lang}_FLAGS "--offload-arch=native" CACHE STRING "Compiler arguments for OpenMP offloading in ${_lang}") + set(TEST_SUITE_OFFLOADING_${_lang}_LDFLAGS "--offload-arch=native" CACHE STRING "Linker arguments for OpenMP offloading in ${_lang}") + if (TARGET OpenMP::OpenMP_${_lang}) + target_link_libraries(OpenMP_Offloading INTERFACE OpenMP::OpenMP_${_lang}) + endif () + separate_arguments(_flags NATIVE_COMMAND "${TEST_SUITE_OFFLOADING_${_lang}_FLAGS}") + separate_arguments(_ldflags NATIVE_COMMAND "${TEST_SUITE_OFFLOADING_${_lang}_LDFLAGS}") + target_compile_options(OpenMP_Offloading INTERFACE $<$:${_flags}>) + target_link_options(OpenMP_Offloading INTERFACE $<$:${_ldflags}>) +endforeach () diff --git a/External/sollve_vv/CMakeLists.txt b/External/sollve_vv/CMakeLists.txt index 1d89a6d0b..0d0f0a43f 100644 --- a/External/sollve_vv/CMakeLists.txt +++ b/External/sollve_vv/CMakeLists.txt @@ -1,138 +1,89 @@ # SOLLVE OpenMP Offloading Validation & Verification Suite # https://crpl.cis.udel.edu/ompvvsollve/ -include(External) - -remove_definitions(-w) -remove_definitions(-Werror=date-time) - -option(TEST_SUITE_FORCE_ALL "Execute all SOLLVE V&V tests, even those known to be unsupported by Clang" OFF) - -set(TEST_SUITE_OFFLOADING_C_FLAGS --offload-arch=native CACHE STRING "Compiler arguments for OpenMP offloading for C") -set(TEST_SUITE_OFFLOADING_CXX_FLAGS --offload-arch=native CACHE STRING "Compiler arguments for OpenMP offloading for CXX") -set(TEST_SUITE_OFFLOADING_Fortran_FLAGS --offload-arch=native CACHE STRING "Compiler arguments for OpenMP offloading for Fortran") -set(TEST_SUITE_OFFLOADING_C_LDFLAGS --offload-arch=native CACHE STRING "Linker arguments for OpenMP offloading") -set(TEST_SUITE_OFFLOADING_CXX_LDFLAGS --offload-arch=native CACHE STRING "Linker arguments for OpenMP offloading") -set(TEST_SUITE_OFFLOADING_Fortran_LDFLAGS --offload-arch=native CACHE STRING "Linker arguments for OpenMP offloading") +if (TEST_SUITE_BENCHMARKING_ONLY) + message(STATUS "Skipping OpenMP Validiation & Verification because benchmarking-only mode") + return () +endif () -set(ALL_LIST_OPTIONS - AMD - INTEL - NVIDIA - REGRESSION -) +include(External) +llvm_externals_find(TEST_SUITE_SOLLVEVV_ROOT "sollve_vv" "OpenMP Offloading Validation & Verification Suite") +if (NOT TEST_SUITE_SOLLVEVV_ROOT) + message(STATUS "Skipping OpenMP Validiation & Verification because external sources not specified") + return () +endif () -set(CHOOSEN_LISTS) +include(../OpenMPOffloading.cmake) +if (NOT TARGET OpenMP_Offloading) + message(STATUS "Skipping OpenMP Validiation & Verification because OpenMP Offloading not available") + return () +endif () -set(REGRESSION) -set(INTEL) +message(STATUS "Adding OpenMP Offloading Validiation & Verification tests") +if (TEST_SUITE_EFFECTIVE_SYSTEM_GPU STREQUAL "all") + file(GLOB_RECURSE _tests_sources RELATIVE "${TEST_SUITE_SOLLVEVV_ROOT}/tests" + "${TEST_SUITE_SOLLVEVV_ROOT}/tests/*.c" + "${TEST_SUITE_SOLLVEVV_ROOT}/tests/*.cpp" + "${TEST_SUITE_SOLLVEVV_ROOT}/tests/*.F90" + ) + set(TEST_SUITE_SOLLVEVV_TESTS ${_tests_sources}) + message(STATUS "Using all SOLLVE V&V tests") +else () + include(sollvevv_${TEST_SUITE_EFFECTIVE_SYSTEM_GPU}_tests.cmake) +endif () -include(sollvevv_nvidia_tests.cmake) -include(sollvevv_amd_tests.cmake) -function (add_sollvevv LANG) +function (add_sollvevv lang) set(_includedir "${TEST_SUITE_SOLLVEVV_ROOT}/ompvv" ) - if ("${LANG}" STREQUAL "Fortran" AND NOT TEST_SUITE_FORTRAN) + if (lang STREQUAL "Fortran" AND NOT TEST_SUITE_FORTRAN) + message(STATUS "Skipping SOLLVE V&V Fortran tests in because Fortran tests are disabled") return () endif () - if (NOT OpenMP_${LANG}_FOUND) - message(FATAL_ERROR "OpenMP for ${LANG} not found") + if (NOT OpenMP_${lang}_FOUND) + message(STATUS "Skipping SOLLVE V&V ${lang} tests in because no OpenMP for ${lang} found") return () endif () - #if (OpenMP_${LANG}_VERSION VERSION_LESS "4.5") - # message(FATAL_ERROR "OpenMP version ${OpenMP_${LANG}_VERSION} too old") - #endif () - - if ("${LANG}" STREQUAL "C") + if (lang STREQUAL "C") set(_langext ".c") - elseif ("${LANG}" STREQUAL "CXX") + elseif (lang STREQUAL "CXX") set(_langext ".cpp") - elseif ("${LANG}" STREQUAL "Fortran") + elseif (lang STREQUAL "Fortran") set(_langext ".F90") else () - message(FATAL_ERROR "Unsupported languge ${LANG}") + return () endif () - file(GLOB_RECURSE _tests_sources RELATIVE "${TEST_SUITE_SOLLVEVV_ROOT}/tests" "${TEST_SUITE_SOLLVEVV_ROOT}/tests/*${_langext}" ) - foreach (_file IN LISTS _tests_sources) + foreach (_file IN LISTS TEST_SUITE_SOLLVEVV_TESTS) get_filename_component(_ext "${_file}" EXT) get_filename_component(_basename "${_file}" NAME_WE) get_filename_component(_directory "${_file}" DIRECTORY) + + if (NOT (_ext STREQUAL _langext)) + continue () + endif () + string(REPLACE "." "" _ext "${_ext}") string(REPLACE "/" "_" _directory "${_directory}") string(REPLACE "." "" _directory "${_directory}") set(_name "omptargetvv-${_basename}-${_directory}-${_ext}") - if (NOT TEST_SUITE_FORCE_ALL AND NOT "${_file}" IN_LIST CHOOSEN_LISTS) - message(STATUS "Skipping SOLLVE V&V test ${_file}") - continue () - endif () - # Create a directory for the test - set(test_dir "${CMAKE_BINARY_DIR}/${_name}") - file(MAKE_DIRECTORY ${test_dir}) + set(_module_dir "${CMAKE_CURRENT_BINARY_DIR}/modules/${_name}") + file(MAKE_DIRECTORY ${_module_dir}) llvm_test_run() llvm_test_executable(${_name} "${TEST_SUITE_SOLLVEVV_ROOT}/tests/${_file}") target_include_directories(${_name} PRIVATE "${_includedir}") - target_link_libraries(${_name} PUBLIC OpenMP::OpenMP_${_lang} m) - - # Add -fopenmp to linker command line; for some reason this is not done by target_link_libraries. - target_link_options(${_name} PRIVATE ${OpenMP_${LANG}_FLAGS}) - - set_target_properties(${_name} PROPERTIES Fortran_MODULE_DIRECTORY ${test_dir}) - - # CMake's find_package(OpenMP) currently does not not introspect flags necessary for offloading. - target_compile_options(${_name} PUBLIC ${TEST_SUITE_OFFLOADING_${LANG}_FLAGS}) - - target_link_options(${_name} PUBLIC ${TEST_SUITE_OFFLOADING_${LANG}_LDFLAGS}) + target_link_libraries(${_name} PUBLIC OpenMP_Offloading m) + set_target_properties(${_name} PROPERTIES Fortran_MODULE_DIRECTORY ${_module_dir}) endforeach () endfunction () - -llvm_externals_find(TEST_SUITE_SOLLVEVV_ROOT "sollve_vv" "OpenMP Offloading Validation & Verification Suite") - -if(TEST_SUITE_SOLLVEVV_ROOT AND NOT TEST_SUITE_BENCHMARKING_ONLY) - if(OpenMP_FOUND) - message(STATUS "Adding OpenMP Offloading Validiation & Verification") - else() - message(STATUS "NOT using OpenMP Validiation & Verification because OpenMP was not found") - return() - endif() - - list(REMOVE_DUPLICATES SYSTEM_GPU) - foreach(list_option ${SYSTEM_GPU}) - string(TOUPPER ${list_option} list_option) - if(list_option IN_LIST ALL_LIST_OPTIONS) - if(list_option STREQUAL "AMD") - list(APPEND CHOOSEN_LISTS ${AMD}) - message(STATUS "adding AMD") - endif() - if(list_option STREQUAL "NVIDIA") - list(APPEND CHOOSEN_LISTS ${NVIDIA}) - message(STATUS "adding NVIDIA") - endif() - if(list_option STREQUAL "INTEL") - list(APPEND CHOOSEN_LISTS ${INTEL}) - message(STATUS "adding INTEL") - endif() - if(list_option STREQUAL "REGRESSION") - list(APPEND CHOOSEN_LISTS ${REGRESSION}) - message(STATUS "adding REGRESSION") - endif() - else() - message(STATUS "Option is unrecognized (${list_option})") - endif() - endforeach() - list(REMOVE_DUPLICATES CHOOSEN_LISTS) - - foreach (_lang in C CXX Fortran) - if(CMAKE_${_lang}_COMPILER) - add_sollvevv(${_lang}) - endif() - endforeach () -endif () +get_property(_enabled_languages GLOBAL PROPERTY ENABLED_LANGUAGES) +foreach (_lang IN LISTS _enabled_languages) + add_sollvevv(${_lang}) +endforeach () diff --git a/External/sollve_vv/README b/External/sollve_vv/README index 7364f3033..f194e0245 100644 --- a/External/sollve_vv/README +++ b/External/sollve_vv/README @@ -19,55 +19,56 @@ OpenMP V&V suite, compile and run them. That is, running llvm-lit machine. Fortran source files won't be run unless specified. Add TEST_SUITE_FORTRAN=ON -and TEST_SUITE_OFFLOADING_Fortran_FLAGS to the CMAKE build. - -To specify the offloading flags, add TEST_SUITE_OFFLOADING__FLAGS -where is C, CXX, or Fortran to the CMAKE build. - -The CMakeLists.txt contains internal greenlists and redlists that the user -must declare which options to use with SYSTEM_GPU. The user is able to -declare multiple lists to use. - -here are all of the options: --amd (this is the greenlist that contains tests that passes) --nvidia (this is the greenlist that contains tests that passes) --intel (this is the greenlist that contains tests that passes) - - -OpenMP support is autodetected by CMake, but clang requires additional -flags to enable offloading. An example run is: - -$ cmake ../llvm-test-suite -GNinja -DCMAKE_BUILD_TYPE=Release \ - -DTEST_SUITE_SOLLVEVV_ROOT=${HOME}/src/sollve_vv \ - -DTEST_SUITE_LIT=${HOME}/build/llvm-project/release/bin/llvm-lit \ - -DCMAKE_C_COMPILER=${HOME}/install/llvm-project/release/bin/clang \ - -DCMAKE_CXX_COMPILER=${HOME}/install/llvm-project/release/bin/clang++ \ - -DTEST_SUITE_SUBDIRS=External/sollve_vv \ - -DTEST_SUITE_SOLLVEVV_OFFLOADING_C_FLAGS=-fopenmp-targets=nvptx64-nvidia-cuda;--cuda-path=/soft/compilers/cuda/cuda-10.1.243;-Xopenmp-target;-march=sm_70 \ - -DTEST_SUITE_SOLLVEVV_OFFLOADING_CXX_FLAGS=-fopenmp-targets=nvptx64-nvidia-cuda;--cuda-path=/soft/compilers/cuda/cuda-10.1.243;-Xopenmp-target;-march=sm_70 \ - -DTEST_SUITE_SOLLVEVV_OFFLOADING_LDFLAGS=-fopenmp-targets=nvptx64-nvidia-cuda;--cuda-path=/soft/compilers/cuda/cuda-10.1.243;-Xopenmp-target;-march=sm_70;-lopenmptarget \ - -DSYSTEM_GPU="amd\;amd_runtime_redlist" \ - -DTEST_SUITE_LIT_FLAGS=-svj1 - -$ LD_LIBRARY_PATH=${HOME}/install/llvm-project/release/lib ninja check - +to the CMake build. + +To specify the offloading flags, add TEST_SUITE_OFFLOADING__FLAGS and +TEST_SUITE_OFFLOADING__LDFLAGS where is C, CXX, or Fortran to the +CMake build. It defaults to `--offload-arch=native`, which is Clang's command +line option to compile for the GPU(s) installed in the current machine. + +Not all features that the SOLLVE V&V suite is testing are currently supported +by Clang/Flang. To avoid tests that are known to always fail, the CMake option +`TEST_SUITE_SYSTEM_GPU` can be used to pre-select a set of tests are are known +to work. Available options are: + + * `-DTEST_SUITE_SYSTEM_GPU=amd` (select tests that should work with Clang/Flang on recent AMD GPUs) + * `-DTEST_SUITE_SYSTEM_GPU=nvidia` (select tests that should work with Clang/Flang on recent Nvidia GPUs) + * `-DTEST_SUITE_SYSTEM_GPU=native` (select one of the options above based on which GPU is installed) + * `-DTEST_SUITE_SYSTEM_GPU=all` (select all SOLLVE V&V tests) + * `-DTEST_SUITE_SYSTEM_GPU=auto` (Use suitable preselected list, otherwise `all`) + + +An example run is: + + $ COMMON_OFFLOAD="-fopenmp-targets=nvptx64-nvidia-cuda --cuda-path=/soft/compilers/cuda/cuda-10.1.243 -Xopenmp-target -march=sm_70" \ + cmake ../llvm-test-suite -GNinja -DCMAKE_BUILD_TYPE=Release \ + -DTEST_SUITE_SOLLVEVV_ROOT=${HOME}/src/sollve_vv \ + -DTEST_SUITE_LIT=${HOME}/build/llvm-project/release/bin/llvm-lit \ + -DCMAKE_C_COMPILER=${HOME}/install/llvm-project/release/bin/clang \ + -DCMAKE_CXX_COMPILER=${HOME}/install/llvm-project/release/bin/clang++ \ + -DTEST_SUITE_SUBDIRS=External/sollve_vv \ + "-DTEST_SUITE_OFFLOADING_C_FLAGS=${COMMON_OFFLOAD}" \ + "-DTEST_SUITE_OFFLOADING_CXX_FLAGS=${COMMON_OFFLOAD}" \ + "-DTEST_SUITE_OFFLOADING_C_LDFLAGS=${COMMON_OFFLOAD}" \ + "-DTEST_SUITE_OFFLOADING_CXX_LDFLAGS=${COMMON_OFFLOAD}" \ + -DTEST_SUITE_SYSTEM_GPU=nvidia \ + -DTEST_SUITE_LIT_FLAGS=-svj1 + + $ LD_LIBRARY_PATH=${HOME}/install/llvm-project/release/lib ninja check Clang also needs to be compiled with enabled offloading for the chosen target. A configuration compatible for the commands above is: -$ cmake ../llvm-project/llvm -GNinja -DCMAKE_BUILD_TYPE=Release \ - -DCMAKE_INSTALL_PREFIX=${HOME}/install/llvm-project/release \ - -DLLVM_ENABLE_PROJECTS=clang;openmp;offload;flang;mlir \ - -DCUDA_TOOLKIT_ROOT_DIR=/soft/compilers/cuda/cuda-10.1.243 \ - -DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES=70 - -$ cmake install + $ cmake ../llvm-project/llvm -GNinja -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX=${HOME}/install/llvm-project/release \ + "-DLLVM_ENABLE_PROJECTS=clang" \ + "-DLLVM_ENABLE_RUNTIMES=openmp;offload" \ + -DCUDA_TOOLKIT_ROOT_DIR=/soft/compilers/cuda/cuda-10.1.243 + $ ninja install In this example, Clang is not installed into a default search path such -that the paths have to be specified explicitly. The options -"-DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES=70" and "-march=sm_70" are to -select the PTX version to compile to. "-DCUDA_TOOLKIT_ROOT_DIR" and -"--cuda-path" point to the CUDA SDK to use. The option +that the paths have to be specified explicitly. "-DCUDA_TOOLKIT_ROOT_DIR" +and "--cuda-path" point to the CUDA SDK to use. The option "-DTEST_SUITE_LIT_FLAGS=-j1" is required to not run the tests in parallel; multiple tests may conflict while running on the same GPU and thus fail. diff --git a/External/sollve_vv/sollvevv_amd_tests.cmake b/External/sollve_vv/sollvevv_amd_tests.cmake index 93e0d226f..08cbf89d7 100644 --- a/External/sollve_vv/sollvevv_amd_tests.cmake +++ b/External/sollve_vv/sollvevv_amd_tests.cmake @@ -1,8 +1,5 @@ -# SOLLVE OpenMP Offloading Validation & Verification Suite -# https://crpl.cis.udel.edu/ompvvsollve/ -# Maintained By: Aaron Jarmusch - -set(AMD +message(STATUS "Using SOLLVE V&V tests known to work with Clang/Flang on AMD GPUs") +set(TEST_SUITE_SOLLVEVV_TESTS 4.5/application_kernels/alpaka_complex_template.cpp 4.5/application_kernels/gemv_target.cpp 4.5/application_kernels/gemv_target_many_matrices.cpp @@ -182,7 +179,6 @@ set(AMD 5.0/master_taskloop/test_master_taskloop.c 5.0/master_taskloop_simd/test_master_taskloop_simd.c 5.0/metadirective/test_metadirective_arch_is_nvidia.c - 5.0/metadirective/test_metadirective_arch_is_nvidia.c 5.0/metadirective/test_metadirective_arch_nvidia_or_amd.c 5.0/parallel_for/test_parallel_for_allocate.c 5.0/parallel_for/test_parallel_for_notequals.c @@ -322,4 +318,4 @@ set(AMD 5.0/requires/test_requires_atomic_default_mem_order_seq_cst.F90 5.0/simd/test_simd_if.F90 4.5/offloading_success.F90 -) # AMD +) diff --git a/External/sollve_vv/sollvevv_nvidia_tests.cmake b/External/sollve_vv/sollvevv_nvidia_tests.cmake index 3e23c25a7..c6da9d6c1 100644 --- a/External/sollve_vv/sollvevv_nvidia_tests.cmake +++ b/External/sollve_vv/sollvevv_nvidia_tests.cmake @@ -1,9 +1,5 @@ -# SOLLVE OpenMP Offloading Validation & Verification Suite -# https://crpl.cis.udel.edu/ompvvsollve/ -# Maintained By: Aaron Jarmusch - - -set(NVIDIA +message(STATUS "Using SOLLVE V&V tests known to work with Clang/Flang on Nvidia GPUs") +set(TEST_SUITE_SOLLVEVV_TESTS 4.5/application_kernels/alpaka_complex_template.cpp 4.5/application_kernels/gemv_target.cpp 4.5/application_kernels/gemv_target_many_matrices.cpp @@ -187,7 +183,6 @@ set(NVIDIA 5.0/master_taskloop_simd/test_master_taskloop_simd.c 5.0/master_taskloop_simd/test_master_taskloop_simd_device.c 5.0/metadirective/test_metadirective_arch_is_nvidia.c - 5.0/metadirective/test_metadirective_arch_is_nvidia.c 5.0/metadirective/test_metadirective_arch_nvidia_or_amd.c 5.0/parallel_for/test_parallel_for_allocate.c 5.0/parallel_for/test_parallel_for_notequals.c @@ -363,4 +358,4 @@ set(NVIDIA 4.5/target_update/test_target_update_from.F90 4.5/target_update/test_target_update_if.F90 4.5/target_update/test_target_update_to.F90 -) # NVIDIA +) From a73e05c7af3f4762cc82d5263a13ece418c35b41 Mon Sep 17 00:00:00 2001 From: jeanPerier Date: Thu, 18 Jul 2024 16:03:30 +0200 Subject: [PATCH 2/9] [fortran] Update gfortran disabled test list now that CLASS is lowered (#109) All the TODOs and bugs related to CLASS have been implemented/fixed. Enable newly passing tests and move all other failing tests under existing or new categories. None of these failures are bugs, most are different TODOs, unimplemented runtime checks, or cases where flang will behave differently than gfortran on purpose. disable class_allocate_19 until further investigation. --- Fortran/gfortran/CMakeLists.txt | 3 + .../gfortran/regression/DisabledFiles.cmake | 78 +++++++++---------- 2 files changed, 39 insertions(+), 42 deletions(-) diff --git a/Fortran/gfortran/CMakeLists.txt b/Fortran/gfortran/CMakeLists.txt index 18cfb2ab4..2d75cdade 100644 --- a/Fortran/gfortran/CMakeLists.txt +++ b/Fortran/gfortran/CMakeLists.txt @@ -138,8 +138,10 @@ set(FLANG_ERRORING_FFLAGS -fdiagnostics-format=json -fdiagnostics-show-option -fdollar-ok + -fdump-fortran-original -fdump-ipa-cp-details -fdump-ipa-fnsummary-details + -fdump-ipa-inline -fdump-ipa-inline-details -fdump-ipa-sra-details -fdump-rtl-combine @@ -149,6 +151,7 @@ set(FLANG_ERRORING_FFLAGS -fdump-tree-cunrolli-details -fdump-tree-dom2 -fdump-tree-dom2-details + -fdump-tree-dse-details -fdump-tree-forwprop2 -fdump-tree-fre1 -fdump-tree-gimple diff --git a/Fortran/gfortran/regression/DisabledFiles.cmake b/Fortran/gfortran/regression/DisabledFiles.cmake index a65aab4e3..734b69c94 100644 --- a/Fortran/gfortran/regression/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/DisabledFiles.cmake @@ -115,6 +115,8 @@ file(GLOB UNSUPPORTED_FILES CONFIGURE_DEPENDS print_fmt_3.f # Unsupported extension intrinsic functions DACOSH, ZCOS, &c. specifics_2.f90 + # Passing Hollerith to CLASS(*). This is not standard. Flang passes it as character. + unlimited_polymorphic_14.f90 # Unsupported predefined macro: __TIMESTAMP__ wdate-time.F90 ) @@ -148,6 +150,8 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS assumed_rank_20.f90 assumed_rank_21.f90 assumed_type_9.f90 + assumed_type_10.f90 + assumed_type_11.f90 bind-c-contiguous-2.f90 interface_49.f90 is_contiguous_2.f90 @@ -160,6 +164,7 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS sizeof_4.f90 sizeof_6.f90 unlimited_polymorphic_1.f03 + unlimited_polymorphic_32.f90 # unimplemented: ASYNCHRONOUS in procedure interface assumed_rank_13.f90 @@ -184,47 +189,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS pr39695_1.f90 pr63797.f90 - # unimplemented: support for polymorphic types. - allocate_alloc_opt_8.f90 - assumed_type_10.f90 - assumed_type_11.f90 - assumed_type_16.f90 - assumed_type_1.f90 - assumed_type_2.f90 - assumed_type_4.f90 - c_f_pointer_tests_5.f90 - class_3.f03 - class_45a.f03 - class_52.f90 - class_allocate_19.f03 - class_array_3.f03 - coarray_allocate_2.f08 - coarray_allocate_3.f08 - coarray_allocate_5.f08 - coarray_lib_alloc_2.f90 - coarray_lib_alloc_3.f90 - coarray_poly_5.f90 - der_io_5.f90 - implicit_class_1.f90 - intent_out_7.f90 - pointer_array_5.f90 - pr105501.f90 - pr48958.f90 - pr57987.f90 - pr63331.f90 - pr78092.f90 - ptr_func_assign_2.f08 - realloc_on_assign_20.f90 - realloc_on_assign_21.f90 - select_type_17.f03 - select_type_48.f90 - sizeof_5.f90 - storage_size_1.f08 - typebound_proc_15.f03 - unlimited_polymorphic_14.f90 - unlimited_polymorphic_3.f03 - unlimited_polymorphic_32.f90 - # unimplemented: BIND (C) internal procedure. ISO_Fortran_binding_19.f90 array_reference_3.f90 @@ -277,12 +241,18 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS coarray_19.f90 coarray_25.f90 coarray_30.f90 + coarray_allocate_2.f08 + coarray_allocate_3.f08 + coarray_allocate_5.f08 coarray_alloc_with_implicit_sync_1.f90 coarray_alloc_with_implicit_sync_2.f90 coarray_lib_alloc_1.f90 + coarray_lib_alloc_2.f90 + coarray_lib_alloc_3.f90 coarray_lib_token_3.f90 coarray_lock_7.f90 coarray_poly_4.f90 + intent_out_7.f90 # unimplemented: coarray in procedure interface coarray_29_1.f90 @@ -294,12 +264,14 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS coarray_lib_move_alloc_1.f90 coarray_lib_this_image_1.f90 coarray_lib_this_image_2.f90 + coarray_poly_5.f90 coarray_poly_6.f90 coarray_poly_7.f90 coarray_poly_8.f90 class_optional_1.f90 class_optional_2.f90 coarray_41.f90 + pr63331.f90 submodule_26.f08 # unimplemented: coarray reference @@ -538,7 +510,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS c_char_tests.f03 # Categorize and debug further - bounds_check_17.f90 # needs -fcheck=bounds c_char_tests_3.f90 # runtime segfault c_funptr_1.f90 # needs "win32_types" module c_funptr_1_mod.f90 # needs "win32_types" module @@ -574,6 +545,10 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS pr69739.f90 typebound_call_32.f90 + # Crashes at compile time in certain builds + # https://github.com/llvm/llvm-project/issues/89179 + class_allocate_19.f03 + # -------------------------------------------------------------------------- # # These tests are skipped because they result in a compile error. This may @@ -1221,6 +1196,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS boz_10.f90 boz_7.f90 byte_1.f90 + class_3.f03 common_19.f90 common_9.f90 constructor_4.f90 @@ -1283,10 +1259,12 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS warnings_are_errors_1.f90 # -std=f2003 tests to detect Fortran 2008 features + allocate_alloc_opt_8.f90 bessel_3.f90 bind_c_array_params.f03 bind_c_bool_1.f90 binding_label_tests_20.f90 + class_52.f90 c_funloc_tests_5.f03 c_kind_int128_test1.f03 c_loc_test_19.f90 @@ -1313,12 +1291,17 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS pointer_target_2.f90 pr95373_2.f90 ptr-func-2.f90 + ptr_func_assign_2.f08 pure_formal_2.f90 rank_2.f90 + realloc_on_assign_20.f90 type_decl_2.f90 typebound_proc_3.f03 + typebound_proc_15.f03 # -std=f2008 tests to detect Fortran 2018 features + assumed_type_16.f90 + assumed_type_4.f90 bind_c_usage_28.f90 c_loc_tests_11.f03 coarray_collectives_2.f90 @@ -1346,6 +1329,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS # Tests that are not errors, and compile just fine with flang constant_shape.f90 continuation_2.f90 + der_io_5.f90 # gfortran's error is inappropriate, array-ctor are not polymorphic. equiv_pure.f90 # gfortran's error is inappropriate for this test finalize_9.f90 func_decl_3.f90 @@ -1354,6 +1338,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS interface_7.f90 intrinsic_short-long.f90 parameter_array_init_6.f90 + pr105501.f90 # Missing space between TYPE and IS is not ambiguous for parser. pr19936_1.f90 # ac-do-variables are their own entities pr35031.f90 # ENTRY in ELEMENTAL pr43996.f90 @@ -1363,6 +1348,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS pr88328.f90 recursive_check_1.f # RECURSIVE is now default recursive_check_2.f90 + select_type_17.f03 # gfortran's error is inappropriate: pointer results are variables. string_1_lp64.f90 substr_10.f90 @@ -1390,6 +1376,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS associated_3.f90 bessel_5_redux.f90 bounds_temporaries_1.f90 + c_f_pointer_tests_5.f90 dec_loc_rval_3.f03 do_4.f do_check_5.f90 @@ -1897,6 +1884,13 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS init_flag_7.f90 init_flag_9.f90 + # Require -fno-realloc-lhs or similar. + realloc_on_assign_21.f90 + + # Require -fcheck=bounds or similar. + bounds_check_17.f90 + pr48958.f90 + # https://github.com/llvm/llvm-project/issues/84088 pr36006-2.f90 ) From 6afc89e737d48753b3618c4d72672e6260e1dac2 Mon Sep 17 00:00:00 2001 From: Kelvin Li Date: Tue, 23 Jul 2024 09:25:12 -0400 Subject: [PATCH 3/9] Check the returns of WIFSIGNALED and WIFEXITED (#144) Need to check WIFSIGNALED and WIFEXTED before using the return values of WTERMSIG and WEXITSTATUS respectively to update variables --- tools/not.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/tools/not.cpp b/tools/not.cpp index ce296b9bb..f45d81f3c 100644 --- a/tools/not.cpp +++ b/tools/not.cpp @@ -92,11 +92,14 @@ int main(int argc, char* const* argv) { retcode = result; signal = 0; } -#elif defined(WEXITSTATUS) && defined(WTERMSIG) +#elif defined(WIFEXITED) && defined(WEXITSTATUS) && defined(WIFSIGNALED) && \ + defined(WTERMSIG) // On POSIX systems and Solaris, result is a composite value of the exit code // and, potentially, the signal that caused termination of the command. - retcode = WEXITSTATUS(result); - signal = WTERMSIG(result); + if (WIFEXITED(result)) + retcode = WEXITSTATUS(result); + if (WIFSIGNALED(result)) + signal = WTERMSIG(result); #else #error "Unsupported system" #endif From 3216058a0cb6a86a3f552bd5b58ec6faa8124b6b Mon Sep 17 00:00:00 2001 From: Tarun Prabhu Date: Thu, 25 Jul 2024 07:24:22 -0600 Subject: [PATCH 4/9] [llvm-test-suite][tools] Fix not utility to avoid std::system (#142) * [llvm-test-suite][tools] Fix not utility to avoid std::system Avoid using std::system on Posix platforms when possible. std::system is now only on non-Posix platforms. When available, posix_spawn is preferred. --- tools/not.cpp | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/tools/not.cpp b/tools/not.cpp index f45d81f3c..5d741a293 100644 --- a/tools/not.cpp +++ b/tools/not.cpp @@ -24,10 +24,9 @@ #include #endif -#ifdef __APPLE__ +#if defined(__unix__) || defined(__APPLE__) #include #include -#include #endif int main(int argc, char* const* argv) { @@ -56,20 +55,22 @@ int main(int argc, char* const* argv) { return 1; int result; -#if !defined(TARGET_OS_IPHONE) + +#if defined(__unix__) || defined(__APPLE__) + pid_t pid; + if (posix_spawn(&pid, argv[0], NULL, NULL, argv, NULL)) + return EXIT_FAILURE; + if (waitpid(pid, &result, WUNTRACED | WCONTINUED) == -1) + return EXIT_FAILURE; +#else std::stringstream ss; ss << argv[0]; for (int i = 1; i < argc; ++i) ss << " " << argv[i]; std::string cmd = ss.str(); result = std::system(cmd.c_str()); -#else - pid_t pid; - if (posix_spawn(&pid, argv[0], NULL, NULL, argv, NULL)) - return EXIT_FAILURE; - if (waitpid(pid, &result, WUNTRACED | WCONTINUED) == -1) - return EXIT_FAILURE; #endif + int retcode = 0; int signal = 0; From 79229479eefcf95bd33c2e3eaa828c33711afda8 Mon Sep 17 00:00:00 2001 From: Tarun Prabhu Date: Thu, 25 Jul 2024 10:51:12 -0600 Subject: [PATCH 5/9] [Fortran/gfortran] Enable tests that now pass Also recategorize some tests that still fail, but for different reasons --- .../gfortran/regression/DisabledFiles.cmake | 188 ++---------------- .../regression/c-interop/DisabledFiles.cmake | 4 - .../regression/ieee/DisabledFiles.cmake | 3 - .../regression/vect/DisabledFiles.cmake | 2 +- 4 files changed, 19 insertions(+), 178 deletions(-) diff --git a/Fortran/gfortran/regression/DisabledFiles.cmake b/Fortran/gfortran/regression/DisabledFiles.cmake index 734b69c94..774a54f98 100644 --- a/Fortran/gfortran/regression/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/DisabledFiles.cmake @@ -130,7 +130,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS ISO_Fortran_binding_1.f90 ISO_Fortran_binding_13.f90 ISO_Fortran_binding_3.f90 - ISO_Fortran_binding_6.f90 PR100029.f90 PR100097.f90 PR100098.f90 @@ -179,9 +178,7 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS recursive_check_6.f03 # unimplemented: procedure components - pr82253.f90 proc_ptr_24.f90 - proc_ptr_comp_14.f90 structure_constructor_11.f90 # unimplemented: procedure pointer results @@ -190,28 +187,21 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS pr63797.f90 # unimplemented: BIND (C) internal procedure. - ISO_Fortran_binding_19.f90 array_reference_3.f90 bind_c_char_4.f90 bind_c_char_5.f90 - internal_dummy_4.f08 # unimplemented: BIND(C) internal procedures: bind-c-char-descr.f90 bind_c_usage_9.f03 # unimplemented: BIND(C) module variable linkage - bind_c_dts.f90 - bind_c_vars.f90 binding_label_tests_10.f03 binding_label_tests_13.f03 global_vars_c_init.f90 proc_ptr_8.f90 # unimplemented: character array expression temp with dynamic length. - char_length_13.f90 - char_result_13.f90 - mapping_3.f90 pr77506.f90 # unimplemented: allocatable components in derived type assignment @@ -227,7 +217,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS # unimplemented: BOZ boz_bge.f90 - nan_4.f90 # unimplemented: coarray address coarray_39.f90 @@ -292,17 +281,9 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS # unimplemented: SYNC MEMORY coarray_sync_memory.f90 - # unimplemented: intrinsic: atan - atan2_1.f90 - # unimplemented: intrinsic: co_broadcast coarray_collectives_17.f90 - # unimplemented: intrinsic: erfc_scaled - erfc_scaled_1.f90 - erf_2.F90 - erf_3.F90 - # Test is not conformant as it expects different value of cmdstat and cmdmsg # Similar test added: UnitTests/execute_command_line execute_command_line_1.f90 @@ -322,10 +303,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS coarray_collectives_12.f90 pr96737.f90 - # unimplemented: intrinsic: selected_char_kind - selected_char_kind_1.f90 - selected_char_kind_4.f90 - # unimplemented: intrinsic: sind dec_math_2.f90 @@ -418,11 +395,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS norm2_3.f90 pr96711.f90 - # unimplemented: calling character elemental function with - # non constant character length (HLFIR regression). - array_temporaries_3.f90 - pure_byref_1.f90 - # unimplemented: pointer assignments inside FORALL (HLFIR # regression) dependency_19.f90 @@ -457,64 +429,31 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # These still crash with flang-new when compiling in this test # environment, but are usually successful when run manually. - ISO_Fortran_binding_8.f90 - ISO_Fortran_binding_12.f90 - binding_c_table_15_1.f03 - binding_label_tests_16.f03 - binding_label_tests_33.f90 - bind_c_coms.f90 bind_c_dts_2.f03 - bind_c_usage_10.f03 - bind_c_usage_16.f03 - bind_c_usage_24.f90 - bind_c_usage_33.f90 - com_block_driver.f90 - c_char_tests_2.f03 c_funloc_tests.f03 c_funloc_tests_3.f03 c_funloc_tests_4.f03 c_f_pointer_shape_tests_5.f90 - c_f_pointer_tests.f90 c_f_pointer_tests_4.f90 - c_kind_int128_test2.f03 c_loc_test.f90 c_loc_tests_2.f03 c_loc_test_20.f90 - c_ptr_tests.f03 c_ptr_tests_14.f90 - c_ptr_tests_19.f90 - c_ptr_tests_7.f03 - c_ptr_tests_8.f03 - c_sizeof_1.f90 - c_sizeof_5.f90 - c_size_t_test.f03 deferred_character_10.f90 iso_c_binding_rename_2.f03 iso_fortran_binding_uint8_array.f90 logical_temp_io.f90 logical_temp_io_kind8.f90 - pointer_remapping_10.f90 - pr32627.f03 pr35983.f90 pr43866.f90 pr71764.f90 pr88611.f90 - print_c_kinds.f90 - proc_decl_17.f90 - proc_decl_2.f90 repack_arrays_1.f90 - transfer_simplify_10.f90 - value_tests_f03.f90 - - # lowering bugs (reported) - c_char_tests.f03 # Categorize and debug further - c_char_tests_3.f90 # runtime segfault c_funptr_1.f90 # needs "win32_types" module c_funptr_1_mod.f90 # needs "win32_types" module c_kind_params.f90 # runtime failure detected in test - class_assign_4.f90 # NYI: vector subscripted polymorphic in HLFIR coarray_allocate_6.f08 # NYI: allocation of coarray coarray_alloc_comp_3.f08 # NYI: allocation of coarray coarray_alloc_comp_6.f08 # NYI: lowering coarray reference @@ -524,20 +463,9 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS coarray_poly_9.f90 # NYI: allocation of coarray c_ptr_tests_10.f03 # valid compilation error on print of c_null_ptr (?) c_ptr_tests_9.f03 # valid compilation error on print of c_null_ptr (?) - ISO_Fortran_binding_10.f90 # test reports failure at runtime - ISO_Fortran_binding_18.f90 # runtime abort - ISO_Fortran_binding_5.f90 # test reports failure at runtime - ISO_Fortran_binding_7.f90 # test reports failure at runtime winapi.f90 # needs -lkernel32 and target *-*-cygwin* widechar_11.f90 # No ASSIGNMENT matches TYPE(c_ptr) and TYPE(__builtin_c_ptr) - # error: 'fir.convert' op invalid type conversion - c_char_tests_4.f90 - c_char_tests_5.f90 - - # error: a function must have a type - proc_decl_9.f90 - # error: not a constant derived type expression coarray_42.f90 init_flag_10.f90 @@ -545,10 +473,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS pr69739.f90 typebound_call_32.f90 - # Crashes at compile time in certain builds - # https://github.com/llvm/llvm-project/issues/89179 - class_allocate_19.f03 - # -------------------------------------------------------------------------- # # These tests are skipped because they result in a compile error. This may @@ -610,7 +534,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # error: '[var]' is an external procedure without the EXTERNAL attribute in # a scope with IMPLICIT NONE(EXTERNAL) - assumed_type_13.f90 bind-c-contiguous-3.f90 # error: Assumed type argument requires an explicit interface @@ -661,11 +584,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS cray_pointers_7.f90 dec_math.f90 dec_math_5.f90 - fmt_en.f90 - fmt_en_rd.f90 - fmt_en_rn.f90 - fmt_en_ru.f90 - fmt_en_rz.f90 fmt_g0_6.f08 fmt_pf.f90 interface_12.f90 @@ -704,12 +622,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # error: 'a' has corank 0, but coindexed reference has 1 cosubscripts coindexed_1.f90 - # error: 'temp_node_t' is PRIVATE in 'temp_node' - constructor_6.f90 - - # error: Type of Cray pointee 'dpte1' is a non-sequence derived type - cray_pointers_2.f90 - # error: DATA statement value could not be converted to the type '[TYPE]' dec_char_conversion_in_data_1.f90 dec_char_conversion_in_data_2.f90 @@ -769,12 +681,7 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS gamma_1.f90 specifics_1.f90 - # error: Keyword may not appear in a reference to a procedure with an implicit - # interface - getenv_1.f90 - # error: Must be a constant value - pr67140.f90 pr89077.f90 substr_simplify.f90 transfer_simplify_12.f90 @@ -797,7 +704,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS pdt_2.f03 # error: '[SYM]' not found in module 'iso_fortran_env' - quad_2.f90 quad_3.f90 team_change_1.f90 team_end_1.f90 @@ -887,13 +793,8 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # Assertion `ty.isa()' failed c_assoc.f90 - c_f_pointer_complex.f03 - c_f_pointer_logical.f03 - c_f_pointer_shape_tests_2.f03 - c_f_pointer_shape_tests_4.f03 equiv_7.f90 iso_c_binding_rename_1.f03 - test_only_clause.f90 # -------------------------------------------------------------------------- # @@ -911,7 +812,6 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # platforms. large_integer_kind.f90 - oldstyle_5.f maxlocval_1.f90 pr91497.f90 @@ -925,16 +825,12 @@ file(GLOB SKIPPED_FILES CONFIGURE_DEPENDS # to do with flang, but they will be skipped until the test suite build # scripts are fixed. - bind_c_array_params_3.f90 class_4a.f03 - global_vars_f90_init.f90 matmul_bounds_14.f namelist_83.f90 pr37287-1.f90 pr77420_3.f90 - pr93524.f90 public_private_module_3.f90 - static_linking_1.f # ---------------------------------------------------------------------------- # @@ -988,12 +884,13 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS bounds_check_strlen_4.f90 bounds_check_strlen_5.f90 bounds_check_strlen_7.f90 + c_char_tests_4.f90 + c_char_tests_5.f90 char_bounds_check_fail_1.f90 char_pointer_assign_4.f90 char_pointer_assign_5.f90 check_bits_1.f90 # requires -fcheck=bits to catch ISHFTC runtime error check_bits_2.f90 # requires -fcheck=bits to catch ISHFTC runtime error - internal_dummy_2.f08 # causes flang-new to crash in the backend (llvm-project/issues/76927) cr_lf.f90 # shenanigans with CR characters dollar_edit_descriptor_4.f # TODO: (i3,$) format shouldn't advance record when looping list_read_11.f90 # more CR character shenanigans @@ -1005,7 +902,6 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS matmul_bounds_8.f90 maxlocval_2.f90 maxlocval_4.f90 - merge_bits_2.F90 merge_char_3.f90 module_nan.f90 namelist_87.f90 @@ -1069,7 +965,6 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS unf_read_corrupted_1.f90 unf_short_record_1.f90 unformatted_subrecord_1.f90 - where_2.f90 widechar_IO_4.f90 zero_sized_1.f90 elemental_function_2.f90 @@ -1084,67 +979,12 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS # # https://github.com/llvm/llvm-test-suite/pull/102#issuecomment-1980674221 # - abort_shouldfail.f90 - all_bounds_1.f90 - alloc_comp_class_4.f03 - allocate_error_5.f90 - allocate_error_6.f90 - allocate_with_source_23.f03 - bounds_check_12.f90 - bounds_check_array_ctor_4.f90 - bounds_check_fail_3.f90 - cshift_bounds_3.f90 - cshift_bounds_4.f90 - deallocate_error_2.f90 - dim_sum_1.f90 - dim_sum_2.f90 - dim_sum_3.f90 - do_check_1.f90 - endfile_3.f90 - eoshift_bounds_1.f90 - error_format.f90 - fmt_error_4.f90 - fmt_error_5.f90 - inline_matmul_15.f90 - internal_write_1.f90 - large_unit_1.f90 - matmul_5.f90 - matmul_bounds_11.f90 - matmul_bounds_13.f90 - matmul_bounds_15.f - matmul_bounds_16.f - matmul_bounds_7.f90 - matmul_bounds_9.f90 - maxloc_bounds_1.f90 - maxloc_bounds_2.f90 - maxloc_bounds_3.f90 - maxloc_bounds_4.f90 - maxloc_bounds_6.f90 - maxloc_bounds_7.f90 - maxloc_bounds_8.f90 - open_errors_2.f90 - open_new_segv.f90 - PR100136.f90 - pack_bounds_1.f90 - pointer_check_11.f90 pointer_check_1.f90 pointer_check_2.f90 pointer_check_3.f90 pointer_check_4.f90 - pointer_check_6.f90 - pr92050.f90 random_3.f90 - random_5.f90 - repeat_1.f90 - reshape_order_1.f90 - reshape_order_2.f90 - reshape_order_3.f90 - reshape_order_4.f90 - spread_bounds_1.f90 unpack_bounds_1.f90 - unpack_bounds_2.f90 - unpack_bounds_3.f90 - write_check.f90 # --------------------------------------------------------------------------- # @@ -1175,7 +1015,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS binding_label_tests_26b.f90 test_common_binding_labels_2_main.f03 - string_1.f90 # Expect error on 32 bits platform + string_1.f90 # Expect error on 32 bits platform # Tests that exercise gfortran's ability to set -std=f95 and then see errors on newer features abstract_type_1.f90 @@ -1425,7 +1265,6 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS interface_6.f90 interop_params.f03 iso_c_binding_class.f03 - pr85877.f90 # Tests that would be errors if we supported options to enable checks dec_structure_24.f90 @@ -1528,8 +1367,8 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS weak-3.f90 # Test is not conformant as it writes to a constant argument # Similar test, that is conformant, added to UnitTests/assign-goto - assign_5.f90 - + assign_5.f90 + # Probable bugs # ["a", "ab"] @@ -1605,8 +1444,6 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS # Valid error: 'x' may not be a local variable in a pure subprogram because: 'x' is polymorphic in a pure subprogram class_49.f90 class_74.f90 - # Valid error: Coindexed polymorphic object may not be associated with a polymorphic dummy argument 'x=' - class_dummy_4.f03 # Valid error: 'atom=' argument must be a scalar coarray or coindexed object for intrinsic 'atomic_ref' coarray_atomic_6.f90 # Valid error: The event-variable must be a coarray @@ -1891,6 +1728,17 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS bounds_check_17.f90 pr48958.f90 - # https://github.com/llvm/llvm-project/issues/84088 - pr36006-2.f90 + # These files require the __truncsfbf2 intrinsic that is not available in + # before GCC 13. Alternatively, it requires compiler-rt to be built and a + # command line option provided to instruct the compiler to use it. Currently, + # we do not support either a version check on GCC or require that compiler-rt + # be built, so these are disabled. See: + # + # https://github.com/llvm/llvm-test-suite/pull/143#discussion_r1689462248 + # + fmt_en.f90 + fmt_en_rd.f90 + fmt_en_rn.f90 + fmt_en_ru.f90 + fmt_en_rz.f90 ) diff --git a/Fortran/gfortran/regression/c-interop/DisabledFiles.cmake b/Fortran/gfortran/regression/c-interop/DisabledFiles.cmake index 903f14e6d..5c9043069 100644 --- a/Fortran/gfortran/regression/c-interop/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/c-interop/DisabledFiles.cmake @@ -38,15 +38,11 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS ubound.f90 # unimplemented: BIND(C) internal procedures - deferred-character-2.f90 fc-out-descriptor-5.f90 - ff-descriptor-5.f90 ff-descriptor-6.f90 - ff-descriptor-7.f90 # unimplemented: support for polymorphic types c407a-1.f90 - c407b-1.f90 ) # These tests are skipped because they trigger internal compiler errors. diff --git a/Fortran/gfortran/regression/ieee/DisabledFiles.cmake b/Fortran/gfortran/regression/ieee/DisabledFiles.cmake index d0fdf1cdc..11454c3db 100644 --- a/Fortran/gfortran/regression/ieee/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/ieee/DisabledFiles.cmake @@ -28,9 +28,6 @@ file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS large_2.f90 rounding_1.f90 rounding_3.f90 - signaling_1.f90 - signaling_2.f90 - signaling_3.f90 # unimplemented: no math runtime available for 'sqrt(f80)' large_1.f90 diff --git a/Fortran/gfortran/regression/vect/DisabledFiles.cmake b/Fortran/gfortran/regression/vect/DisabledFiles.cmake index 94bb490d1..4f0dc3c44 100644 --- a/Fortran/gfortran/regression/vect/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/vect/DisabledFiles.cmake @@ -1,3 +1,4 @@ + #===------------------------------------------------------------------------===# # # Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -21,6 +22,5 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS pr90681.f pr97761.f90 pr99746.f90 - vect-8.f90 vect-8-epilogue.F90 ) From 4a138135513ca5d6e03abcfa11cd9f1576442b4d Mon Sep 17 00:00:00 2001 From: Tarun Prabhu Date: Thu, 25 Jul 2024 10:53:01 -0600 Subject: [PATCH 6/9] [Fortran/gfortran] Remove HLFIR-specific allow and deny lists The code that supports "feature-specific" testing has been retained, but has been tweaked a little bit. The README has been updated with some examples. --- Fortran/gfortran/CMakeLists.txt | 21 +++---- Fortran/gfortran/README.md | 33 +++++++++-- .../regression/DisabledFilesHLFIR.cmake | 59 ------------------- .../regression/EnabledFilesHLFIR.cmake | 37 ------------ .../analyzer/DisabledFilesHLFIR.cmake | 7 --- .../analyzer/EnabledFilesHLFIR.cmake | 7 --- .../regression/asan/DisabledFilesHLFIR.cmake | 7 --- .../regression/asan/EnabledFilesHLFIR.cmake | 7 --- .../c-interop/DisabledFilesHLFIR.cmake | 7 --- .../c-interop/EnabledFilesHLFIR.cmake | 7 --- .../coarray/DisabledFilesHLFIR.cmake | 7 --- .../coarray/EnabledFilesHLFIR.cmake | 7 --- .../regression/debug/DisabledFilesHLFIR.cmake | 7 --- .../regression/debug/EnabledFilesHLFIR.cmake | 7 --- .../regression/g77/DisabledFilesHLFIR.cmake | 7 --- .../regression/g77/EnabledFilesHLFIR.cmake | 7 --- .../goacc-gomp/DisabledFilesHLFIR.cmake | 7 --- .../goacc-gomp/EnabledFilesHLFIR.cmake | 7 --- .../regression/goacc/DisabledFilesHLFIR.cmake | 7 --- .../regression/goacc/EnabledFilesHLFIR.cmake | 7 --- .../regression/gomp/DisabledFilesHLFIR.cmake | 7 --- .../regression/gomp/EnabledFilesHLFIR.cmake | 7 --- .../gomp/appendix-a/DisabledFilesHLFIR.cmake | 7 --- .../gomp/appendix-a/EnabledFilesHLFIR.cmake | 7 --- .../regression/ieee/DisabledFilesHLFIR.cmake | 7 --- .../regression/ieee/EnabledFilesHLFIR.cmake | 7 --- .../regression/lto/DisabledFilesHLFIR.cmake | 7 --- .../regression/lto/EnabledFilesHLFIR.cmake | 7 --- .../regression/ubsan/DisabledFilesHLFIR.cmake | 16 ----- .../regression/ubsan/EnabledFilesHLFIR.cmake | 7 --- .../regression/vect/DisabledFilesHLFIR.cmake | 7 --- .../regression/vect/EnabledFilesHLFIR.cmake | 7 --- .../torture/compile/DisabledFilesHLFIR.cmake | 7 --- .../torture/compile/EnabledFilesHLFIR.cmake | 7 --- .../torture/execute/DisabledFilesHLFIR.cmake | 12 ---- .../torture/execute/EnabledFilesHLFIR.cmake | 7 --- 36 files changed, 36 insertions(+), 352 deletions(-) delete mode 100644 Fortran/gfortran/regression/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/analyzer/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/analyzer/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/asan/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/asan/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/c-interop/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/c-interop/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/coarray/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/coarray/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/debug/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/debug/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/g77/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/g77/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/goacc-gomp/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/goacc-gomp/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/goacc/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/goacc/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/gomp/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/gomp/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/gomp/appendix-a/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/gomp/appendix-a/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/ieee/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/ieee/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/lto/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/lto/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/ubsan/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/ubsan/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/vect/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/regression/vect/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/torture/compile/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/torture/compile/EnabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/torture/execute/DisabledFilesHLFIR.cmake delete mode 100644 Fortran/gfortran/torture/execute/EnabledFilesHLFIR.cmake diff --git a/Fortran/gfortran/CMakeLists.txt b/Fortran/gfortran/CMakeLists.txt index 2d75cdade..2ae7b67ed 100644 --- a/Fortran/gfortran/CMakeLists.txt +++ b/Fortran/gfortran/CMakeLists.txt @@ -473,28 +473,25 @@ function(gfortran_populate_disabled_tests out) list(APPEND skipped ${SKIPPED_FILES}) list(APPEND failing ${FAILING_FILES}) - # do the same for any requested feature extentions + # do the same for any requested feature extensions foreach(feature ${TEST_SUITE_FORTRAN_FEATURES}) set(UNSUPPORTED_FILES "") set(UNIMPLEMENTED_FILES "") set(SKIPPED_FILES "") set(FAILING_FILES "") - include(${CMAKE_CURRENT_SOURCE_DIR}/DisabledFiles${feature}.cmake) + include(${CMAKE_CURRENT_SOURCE_DIR}/DisabledFiles_${feature}.cmake) list(APPEND unsupported ${UNSUPPORTED_FILES}) list(APPEND unimplemented ${UNIMPLEMENTED_FILES}) list(APPEND skipped ${SKIPPED_FILES}) list(APPEND failing ${FAILING_FILES}) # enable any tests that now pass for this feature - set(SUPPORTED_FILES "") - set(IMPLEMENTED_FILES "") - set(UNSKIPPED_FILES "") - set(PASSING_FILES "") - include(${CMAKE_CURRENT_SOURCE_DIR}/EnabledFiles${feature}.cmake) - list(REMOVE_ITEM unsupported ${SUPPORTED_FILES}) - list(REMOVE_ITEM unimplemented ${IMPLEMENTED_FILES}) - list(REMOVE_ITEM skipped ${UNSKIPPED_FILES}) - list(REMOVE_ITEM failing ${PASSING_FILES}) + set(ENABLED_FILES "") + include(${CMAKE_CURRENT_SOURCE_DIR}/EnabledFiles_${feature}.cmake) + list(REMOVE_ITEM unsupported ${ENABLED_FILES}) + list(REMOVE_ITEM unimplemented ${ENABLED_FILES}) + list(REMOVE_ITEM skipped ${ENABLED_FILES}) + list(REMOVE_ITEM failing ${ENABLED_FILES}) endforeach() set(disabled "") @@ -511,7 +508,7 @@ function(gfortran_populate_disabled_tests out) # and unimplemented tests since some could be enabled once some feature is # implemented. Eventually, all the TEST_SUITE_FORTRAN_FORCE_* options (perhaps # with the exception of TEST_SUITE_FORTRAN_FORCE_UNSUPPORTED_TESTS) should - # become redundant and removed. + # become redundant and should be removed. if (NOT TEST_SUITE_FORTRAN_FORCE_ALL_TESTS AND NOT TEST_SUITE_FORTRAN_FORCE_UNIMPLEMENTED_TESTS) list(APPEND disabled ${unimplemented}) diff --git a/Fortran/gfortran/README.md b/Fortran/gfortran/README.md index 5006cd978..8076e820b 100644 --- a/Fortran/gfortran/README.md +++ b/Fortran/gfortran/README.md @@ -115,14 +115,35 @@ are unrelated to the gfortran tests here. ### Testing non-standard features/flags Additional denylists for a particular feature can be included by creating -DisabledFilesFEATURE.cmake files (in the same format as those for the default -denylists), and adding FEATURE to `TEST_SUITE_FORTRAN_FEATURES`. Additional -compiler flags can be added using `CMAKE_Fortran_FLAGS`. +`DisabledFiles_FEATURE.cmake` files (in the same format as those for the default +denylists, `DisabledFiles.cmake`), and adding FEATURE to +`TEST_SUITE_FORTRAN_FEATURES`. Additional compiler flags can be added using +`CMAKE_Fortran_FLAGS`. -For example, to test HLFIR one could use -`CMAKE_Fortran_Flags=-flang-experimental-hlfir` and -`TEST_SUITE_FORTRAN_FEATURES=HLFIR`. +For example, to test feature, FOO, one could use +``` +cmake -DTEST_SUITE_FORTRAN_FEATURES=FOO \ + -DCMAKE_Fortran_FLAGS=-some-foo-specific-flag-if-required \ + +``` + +`DisabledFiles_FOO.cmake` files can be created in the appropriate subdirectories +if enabling the feature/flag results in the failure of tests that otherwise pass. +Conversely, the feature/flag may cause some disabled tests to pass. These can be +added to an allowlist file, `EnabledFiles_FOO.cmake` in the corresponding +directory. The file must contain a single variable named `ENABLED_FILES` with +the file names of the tests that should be enabled (in the case of multi-file +tests, this should be the name of the "main" file). An example of such a list is +below. + +``` +file(GLOB ENABLED_FILES CONFIGURE_DEPENDS + test_1.f90 + multifile-main.f03 +) + +``` ### Notes for developers/maintainers ### diff --git a/Fortran/gfortran/regression/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/DisabledFilesHLFIR.cmake deleted file mode 100644 index ab9f33bd4..000000000 --- a/Fortran/gfortran/regression/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,59 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# - -# These tests are skipped because they hit a 'not yet implemented' assertion -# in flang and thus fail to compile. They should be removed from here when the -# corresponding feature is implemented. Eventually, this variable should be -# removed altogether once all the missing features are implemented. -file(GLOB UNIMPLEMENTED_FILES CONFIGURE_DEPENDS - # unimplemented: lower descriptor designator to HLFIR value - allocatable_scalar_1.f90 - deferred_type_param_8.f90 - pointer_function_actual_1.f90 - - # unimplemented: compute elemental function result length parameters in HLFIR - array_temporaries_3.f90 - pure_byref_1.f90 - - # unimplemented: BIND(C) INTENT(OUT) allocatable deallocation in HLFIR - bind_c_procs_3.f90 - - # unimplemented: pointer assignment inside FORALL - dependency_19.f90 - forall_3.f90 - pr49698.f90 -) - -# These tests are disabled because they fail when they are expected to pass. -file(GLOB FAILING_FILES CONFIGURE_DEPENDS - # --------------------------------------------------------------------------- - # - # These tests fail at runtime when they should pass. are likely a result of - # unimplemented features in the runtime, but they could also be bugs. If any - # will never pass with flang (if they use unsupported extensions for instance), - # they should be added to the Unsupported list. - - all_bounds_1.f90 - array_constructor_11.f90 - array_constructor_45.f90 - bounds_check_array_ctor_4.f90 - bounds_check_fail_3.f90 - bounds_check_12.f90 - function_optimize_11.f90 - maxloc_bounds_1.f90 - maxloc_bounds_2.f90 - maxloc_bounds_4.f90 - maxloc_bounds_5.f90 - maxloc_bounds_7.f90 - maxloc_bounds_8.f90 - pack_bounds_1.f90 - pr67524.f90 - spread_bounds_1.f90 - transfer_array_intrinsic_4.f90 - maxlocval_1.f90 -) diff --git a/Fortran/gfortran/regression/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/EnabledFilesHLFIR.cmake deleted file mode 100644 index 1891b03bc..000000000 --- a/Fortran/gfortran/regression/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,37 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# - -# These are files which are skipped by default (see DisabledFiles.cmake) but -# are fixed by this feature -file(GLOB UNSKIPPED_FILES CONFIGURE_DEPENDS - char_cast_2.f90 - deferred_character_8.f90 - dependency_23.f90 - interface_assignment_1.f90 - typebound_assignment_6.f90 - volatile10.f90 -) - -# These are files which fail by default (see DisabledFiles.cmake) but are fixed -# by this feature -file(GLOB PASSING_FILES CONFIGURE_DEPENDS - advance_5.f90 - aliasing_dummy_5.f90 - dependency_45.f90 - elemental_dependency_1.f90 - elemental_dependency_5.f90 - elemental_dependency_6.f90 - forall_12.f90 - forall_17.f90 - inline_transpose_1.f90 - internal_pack_3.f90 - missing_optional_dummy_6.f90 - mvbits_4.f90 - pr50069_1.f90 - pr68227.f90 -) diff --git a/Fortran/gfortran/regression/analyzer/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/analyzer/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/analyzer/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/analyzer/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/analyzer/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/analyzer/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/asan/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/asan/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/asan/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/asan/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/asan/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/asan/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/c-interop/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/c-interop/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/c-interop/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/c-interop/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/c-interop/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/c-interop/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/coarray/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/coarray/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/coarray/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/coarray/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/coarray/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/coarray/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/debug/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/debug/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/debug/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/debug/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/debug/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/debug/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/g77/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/g77/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/g77/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/g77/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/g77/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/g77/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/goacc-gomp/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/goacc-gomp/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/goacc-gomp/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/goacc-gomp/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/goacc-gomp/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/goacc-gomp/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/goacc/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/goacc/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/goacc/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/goacc/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/goacc/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/goacc/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/gomp/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/gomp/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/gomp/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/gomp/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/gomp/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/gomp/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/gomp/appendix-a/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/gomp/appendix-a/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/gomp/appendix-a/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/gomp/appendix-a/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/gomp/appendix-a/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/gomp/appendix-a/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/ieee/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/ieee/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/ieee/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/ieee/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/ieee/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/ieee/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/lto/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/lto/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/lto/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/lto/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/lto/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/lto/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/ubsan/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/ubsan/DisabledFilesHLFIR.cmake deleted file mode 100644 index 8158d7775..000000000 --- a/Fortran/gfortran/regression/ubsan/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,16 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# - -# These tests are skipped because they hit a 'not yet implemented' assertion -# in flang and thus fail to compile. They should be removed from here when the -# corresponding feature is implemented. Eventually, this vairable should be -# removed altogether once all the missing features are implemented. -file (GLOB UNIMPLEMENTED_FILES - # unimplemented: BIND(C) INTENT(OUT) allocatable deallocation in HLFIR - bind-c-intent-out-2.f90 -) diff --git a/Fortran/gfortran/regression/ubsan/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/ubsan/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/ubsan/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/vect/DisabledFilesHLFIR.cmake b/Fortran/gfortran/regression/vect/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/vect/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/regression/vect/EnabledFilesHLFIR.cmake b/Fortran/gfortran/regression/vect/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/regression/vect/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/torture/compile/DisabledFilesHLFIR.cmake b/Fortran/gfortran/torture/compile/DisabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/torture/compile/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/torture/compile/EnabledFilesHLFIR.cmake b/Fortran/gfortran/torture/compile/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/torture/compile/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# diff --git a/Fortran/gfortran/torture/execute/DisabledFilesHLFIR.cmake b/Fortran/gfortran/torture/execute/DisabledFilesHLFIR.cmake deleted file mode 100644 index 395d77fb9..000000000 --- a/Fortran/gfortran/torture/execute/DisabledFilesHLFIR.cmake +++ /dev/null @@ -1,12 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# - -# These tests are disabled because they fail at runtime when they should pass. -file(GLOB FAILING_FILES CONFIGURE_DEPENDS - constructor.f90 -) diff --git a/Fortran/gfortran/torture/execute/EnabledFilesHLFIR.cmake b/Fortran/gfortran/torture/execute/EnabledFilesHLFIR.cmake deleted file mode 100644 index 2bb24654e..000000000 --- a/Fortran/gfortran/torture/execute/EnabledFilesHLFIR.cmake +++ /dev/null @@ -1,7 +0,0 @@ -#===------------------------------------------------------------------------===# -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -#===------------------------------------------------------------------------===# From 00e2fa66423c7c155c2f5cb797ad9c3ea7695626 Mon Sep 17 00:00:00 2001 From: Fabian Ritter Date: Mon, 29 Jul 2024 14:56:51 +0200 Subject: [PATCH 7/9] [HIP] Add a correctness test for the memmove intrinsic (#146) So far, the lowering of device-side memmove intrinsics in HIP (and the AMDGPU backend) is only tested with syntactic regression tests. This patch adds functional correctness tests for the device-side memmove intrinsic with and without overlapping source and destination ranges. By testing various statically known or unknown move lengths, the different lowering mechanisms for memmove in the AMDGPU backend are covered. --- External/HIP/CMakeLists.txt | 1 + External/HIP/memmove.hip | 313 ++++++++++++++++++++++++++ External/HIP/memmove.reference_output | 2 + 3 files changed, 316 insertions(+) create mode 100644 External/HIP/memmove.hip create mode 100644 External/HIP/memmove.reference_output diff --git a/External/HIP/CMakeLists.txt b/External/HIP/CMakeLists.txt index d11add036..c5d1c80fe 100644 --- a/External/HIP/CMakeLists.txt +++ b/External/HIP/CMakeLists.txt @@ -15,6 +15,7 @@ macro(create_local_hip_tests VariantSuffix) list(APPEND HIP_LOCAL_TESTS empty) list(APPEND HIP_LOCAL_TESTS with-fopenmp) list(APPEND HIP_LOCAL_TESTS saxpy) + list(APPEND HIP_LOCAL_TESTS memmove) list(APPEND HIP_LOCAL_TESTS InOneWeekend) list(APPEND HIP_LOCAL_TESTS TheNextWeek) diff --git a/External/HIP/memmove.hip b/External/HIP/memmove.hip new file mode 100644 index 000000000..31d41af8f --- /dev/null +++ b/External/HIP/memmove.hip @@ -0,0 +1,313 @@ +#include +#include +#include +#include + +#include "hip/hip_runtime.h" + +// Tests for the functional correctness of the lowering of memmove in device +// code, including moves with overlapping source and destination ranges. Various +// memmoves are performed on device side and the result of each is compared to +// the corresponding operation on the host. +// Global, shared, and stack memory is tested. + +#define VERBOSE 0 + +#define CHKHIP(r) \ + if (r != hipSuccess) { \ + std::cerr << hipGetErrorString(r) << std::endl; \ + abort(); \ + } + +using item_type = uint8_t; + +// Maximal number of bytes to copy with a memmove call, used to allocate +// buffers. +#define MAX_BYTES_PER_THREAD 2048 + +// LDS is small, so run only smaller tests there. +#define MAX_BYTES_PER_THREAD_SHARED 128 + +// Number of threads that move started in parallel. +#define NUM_MOVE_THREADS (2 * 32) + +// Size of blocks in the grid used for move threads. If the number of threads is +// smaller than this, it is used instead. +#define BLOCK_SIZE 256 + +#define ALLOC_SIZE (2 * NUM_MOVE_THREADS * MAX_BYTES_PER_THREAD) + +#define ALLOC_SIZE_SHARED (2 * NUM_MOVE_THREADS * MAX_BYTES_PER_THREAD_SHARED) + +#define TESTED_FUNCTION __builtin_memmove + +enum AddressSpace { + GLOBAL = 0, + SHARED = 1, + STACK = 2, +}; + +static const char *as_names[] = { + "global", + "shared", + "stack", +}; + +static constexpr size_t get_stride(size_t bytes_per_thread) { + return 2 * bytes_per_thread; +} + +__global__ void init_kernel(item_type *buf_device, size_t alloc_size) { + for (size_t i = 0; i < alloc_size; ++i) { + buf_device[i] = (item_type)i; + } +} + +template +__global__ void move_kernel_global_const(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + (void)dyn_sz; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + item_type *thread_buf = buf_device + get_stride(SZ) * tid; + TESTED_FUNCTION(thread_buf + dst_idx, thread_buf + src_idx, SZ); +} + +template +__global__ void move_kernel_shared_const(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + (void)dyn_sz; + __shared__ item_type buf_shared[ALLOC_SIZE_SHARED]; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + constexpr size_t stride = get_stride(SZ); + item_type *thread_buf = buf_device + stride * tid; + item_type *thread_buf_shared = buf_shared + stride * tid; + // Copy the original data to shared memory. + __builtin_memcpy(thread_buf_shared, thread_buf, stride); + // Perform the move there. + TESTED_FUNCTION(thread_buf_shared + dst_idx, thread_buf_shared + src_idx, SZ); + // Copy the modified data back to global memory. + __builtin_memcpy(thread_buf, thread_buf_shared, stride); +} + +template +__global__ void move_kernel_stack_const(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + (void)dyn_sz; + constexpr size_t stride = get_stride(SZ); + item_type buf_stack[stride]; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + item_type *thread_buf = buf_device + stride * tid; + // Copy the original data to the stack. + __builtin_memcpy(buf_stack, thread_buf, stride); + // Perform the move there. + TESTED_FUNCTION(buf_stack + dst_idx, buf_stack + src_idx, SZ); + // Copy the modified data back to global memory. + __builtin_memcpy(thread_buf, buf_stack, stride); +} + +__global__ void move_kernel_global_var(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + item_type *thread_buf = buf_device + get_stride(dyn_sz) * tid; + TESTED_FUNCTION(thread_buf + dst_idx, thread_buf + src_idx, dyn_sz); +} + +__global__ void move_kernel_shared_var(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + __shared__ item_type buf_shared[ALLOC_SIZE_SHARED]; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + size_t stride = get_stride(dyn_sz); + item_type *thread_buf = buf_device + stride * tid; + item_type *thread_buf_shared = buf_shared + stride * tid; + // Copy the original data to shared memory. + __builtin_memcpy(thread_buf_shared, thread_buf, stride); + // perform the move there + TESTED_FUNCTION(thread_buf_shared + dst_idx, thread_buf_shared + src_idx, + dyn_sz); + // Copy the modified data back to global memory. + __builtin_memcpy(thread_buf, thread_buf_shared, stride); +} + +template +__global__ void move_kernel_stack_var(item_type *buf_device, size_t src_idx, + size_t dst_idx, size_t dyn_sz) { + // We use the static SZ to allocate a fixed-size stack variable. + constexpr size_t stride = get_stride(SZ); + item_type buf_stack[stride]; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + if (tid >= NUM_MOVE_THREADS) + return; + item_type *thread_buf = buf_device + stride * tid; + // Copy the original data to the stack. + __builtin_memcpy(buf_stack, thread_buf, stride); + // perform the move there + TESTED_FUNCTION(buf_stack + dst_idx, buf_stack + src_idx, dyn_sz); + // Copy the modified data back to global memory. + __builtin_memcpy(thread_buf, buf_stack, stride); +} + +template +bool run_test(item_type *buf_reference, item_type *buf_host, + item_type *buf_device, size_t src_idx, size_t dst_idx, + bool const_size, AddressSpace AS, size_t &differing_pos) { + // Initialize device buffer. + hipLaunchKernelGGL(init_kernel, dim3(1), dim3(1), 0, 0, buf_device, + ALLOC_SIZE); + CHKHIP(hipDeviceSynchronize()); + + // Set up the reference buffer. + for (size_t i = 0; i < ALLOC_SIZE; ++i) + buf_reference[i] = (item_type)i; + + // Simulate multi-threaded device-side memmove on the host. + for (size_t tid = 0; tid < NUM_MOVE_THREADS; ++tid) { + item_type *thread_buf = buf_reference + get_stride(SZ) * tid; + std::memmove(thread_buf + dst_idx, thread_buf + src_idx, SZ); + } + + // Do the device-side memmove. + int block_size = std::min(BLOCK_SIZE, NUM_MOVE_THREADS); + int num_blocks = (NUM_MOVE_THREADS + block_size - 1) / block_size; + + switch (AS) { + case AddressSpace::GLOBAL: + hipLaunchKernelGGL(const_size ? move_kernel_global_const + : move_kernel_global_var, + dim3(num_blocks), dim3(block_size), 0, 0, buf_device, + src_idx, dst_idx, SZ); + break; + case AddressSpace::SHARED: + hipLaunchKernelGGL(const_size ? move_kernel_shared_const + : move_kernel_shared_var, + dim3(num_blocks), dim3(block_size), 0, 0, buf_device, + src_idx, dst_idx, SZ); + break; + case AddressSpace::STACK: + hipLaunchKernelGGL(const_size ? move_kernel_stack_const + : move_kernel_stack_var, + dim3(num_blocks), dim3(block_size), 0, 0, buf_device, + src_idx, dst_idx, SZ); + break; + }; + CHKHIP(hipDeviceSynchronize()); + + // Fetch the result into buf_host. + CHKHIP(hipMemcpy(buf_host, buf_device, ALLOC_SIZE, hipMemcpyDeviceToHost)); + + // Compare to the reference. + bool success = true; + for (size_t i = 0; i < ALLOC_SIZE; ++i) { + if (buf_host[i] != buf_reference[i]) { + differing_pos = i; + success = false; + break; + } + } + + return success; +} + +template +int run_tests(item_type *buf_reference, item_type *buf_host, + item_type *buf_device, AddressSpace AS) { + if (AS == AddressSpace::SHARED && SZ > MAX_BYTES_PER_THREAD_SHARED) { + // LDS is too small for these tests. + return 0; + } + assert(SZ <= MAX_BYTES_PER_THREAD && + "Increase MAX_BYTES_PER_THREAD for larger sizes"); + + std::vector> index_combinations = { + {0, 1}, {0, SZ}, {0, SZ - 1}, {1, 0}, {SZ, 0}, {SZ - 1, 0}, + }; + if (SZ > 16) { + index_combinations.emplace_back(0, 16); + index_combinations.emplace_back(16, 0); + } + + int nerrs = 0; + + size_t differing_pos = 0; + auto test_index_combinations = [&](bool const_size) { + for (const auto &[src_idx, dst_idx] : index_combinations) { + bool success = run_test(buf_reference, buf_host, buf_device, src_idx, + dst_idx, const_size, AS, differing_pos); + nerrs += !success; + if (VERBOSE || !success) { + std::cout << "- moving [" << src_idx << ", " << (src_idx + SZ - 1) + << "] -> [" << dst_idx << ", " << (dst_idx + SZ - 1) << "]"; + if (!VERBOSE) { + std::cout << " with " << (const_size ? "static" : "dynamic") + << " size in " << as_names[AS] << " memory"; + } + std::cout << ":"; + if (success) { + std::cout << " successful\n"; + } else { + std::cout << " failed\n -> first difference at index " + << differing_pos << '\n'; + } + } + } + }; + + if (VERBOSE) + std::cout << "running tests for dynamic move length " << SZ << " in " + << as_names[AS] << " memory\n"; + test_index_combinations(false); + + // Different paths in codegen are taken if the move length is statically + // known. + if (VERBOSE) + std::cout << "running tests for static move length " << SZ << " in " + << as_names[AS] << " memory\n"; + test_index_combinations(true); + + return nerrs; +} + +int main(void) { + item_type *buf_device; + CHKHIP(hipMalloc(&buf_device, ALLOC_SIZE)); + + std::unique_ptr buf_host(new item_type[ALLOC_SIZE]); + std::unique_ptr buf_reference(new item_type[ALLOC_SIZE]); + + int nerrs = 0; + for (AddressSpace AS : + {AddressSpace::GLOBAL, AddressSpace::SHARED, AddressSpace::STACK}) { + nerrs += run_tests<64>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += run_tests<66>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += run_tests<73>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += run_tests<3>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += run_tests<1>(buf_reference.get(), buf_host.get(), buf_device, AS); + + // Move lengths that are large enough for the IR lowering in the constant + // case, with simple residual, no residual, and maximal residual: + nerrs += + run_tests<1025>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += + run_tests<1040>(buf_reference.get(), buf_host.get(), buf_device, AS); + nerrs += + run_tests<1039>(buf_reference.get(), buf_host.get(), buf_device, AS); + } + + CHKHIP(hipFree(buf_device)); + + if (nerrs != 0) { + std::cout << nerrs << " errors\n"; + return 1; + } + std::cout << "PASSED!\n"; + return 0; +} diff --git a/External/HIP/memmove.reference_output b/External/HIP/memmove.reference_output new file mode 100644 index 000000000..391efdf64 --- /dev/null +++ b/External/HIP/memmove.reference_output @@ -0,0 +1,2 @@ +PASSED! +exit 0 From 7517cb8c45a31a03ba2f3aa3fa3c703de441d195 Mon Sep 17 00:00:00 2001 From: Tarun Prabhu Date: Fri, 2 Aug 2024 01:52:00 -0600 Subject: [PATCH 8/9] [Fortran/gfortran] Disable test to fix failing buildbots (#149) Several of the AArch64 buildbots are failing due to an undefined reference to __trampoline_setup when compiling internal_dummy_2.f08. This is probably an issue with the buildbot configuration, but in the interest of getting them bots unstuck, the test is being disabled. Co-authored-by: Tarun Prabhu --- Fortran/gfortran/regression/DisabledFiles.cmake | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/Fortran/gfortran/regression/DisabledFiles.cmake b/Fortran/gfortran/regression/DisabledFiles.cmake index 774a54f98..a4e1d5fa0 100644 --- a/Fortran/gfortran/regression/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/DisabledFiles.cmake @@ -1741,4 +1741,9 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS fmt_en_rn.f90 fmt_en_ru.f90 fmt_en_rz.f90 + + # These test causes failures in some buildbots with an undefined reference to + # __trampoline_setup. This is probably an unrelated issue, but as a quick fix + # for the buildbot, this is disabled. + internal_dummy_2.f08 ) From 5e819e187f5f9e91aaeffdb373631700b50d6624 Mon Sep 17 00:00:00 2001 From: antoine moynault Date: Fri, 2 Aug 2024 11:23:08 +0200 Subject: [PATCH 9/9] [Fortran/gfortran] Disable another test to fix failing buildbots (#150) Fortran/gfortran/regression/gfortran-regression-execute-regression__do_check_1_f90.test still fails after being re-enabled in #143 https://lab.llvm.org/buildbot/#/builders/143/builds/1198 --- Fortran/gfortran/regression/DisabledFiles.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/Fortran/gfortran/regression/DisabledFiles.cmake b/Fortran/gfortran/regression/DisabledFiles.cmake index a4e1d5fa0..2a6414089 100644 --- a/Fortran/gfortran/regression/DisabledFiles.cmake +++ b/Fortran/gfortran/regression/DisabledFiles.cmake @@ -979,6 +979,7 @@ file(GLOB FAILING_FILES CONFIGURE_DEPENDS # # https://github.com/llvm/llvm-test-suite/pull/102#issuecomment-1980674221 # + do_check_1.f90 pointer_check_1.f90 pointer_check_2.f90 pointer_check_3.f90