diff --git a/libc/cmake/modules/LLVMLibCCheckMPFR.cmake b/libc/cmake/modules/LLVMLibCCheckMPFR.cmake index 9e361f5fd8112..46f679f1330d3 100644 --- a/libc/cmake/modules/LLVMLibCCheckMPFR.cmake +++ b/libc/cmake/modules/LLVMLibCCheckMPFR.cmake @@ -2,8 +2,6 @@ set(LLVM_LIBC_MPFR_INSTALL_PATH "" CACHE PATH "Path to where MPFR is installed ( if(LLVM_LIBC_MPFR_INSTALL_PATH) set(LIBC_TESTS_CAN_USE_MPFR TRUE) -elseif(LIBC_TARGET_ARCHITECTURE_IS_GPU) - set(LIBC_TESTS_CAN_USE_MPFR FALSE) else() try_compile( LIBC_TESTS_CAN_USE_MPFR diff --git a/libc/test/src/CMakeLists.txt b/libc/test/src/CMakeLists.txt index c45b94f364397..10ad7ff595554 100644 --- a/libc/test/src/CMakeLists.txt +++ b/libc/test/src/CMakeLists.txt @@ -8,7 +8,7 @@ function(add_fp_unittest name) ) if(MATH_UNITTEST_NEED_MPFR) - if(NOT LIBC_TESTS_CAN_USE_MPFR) + if(NOT LIBC_TESTS_CAN_USE_MPFR OR LIBC_TARGET_ARCHITECTURE_IS_GPU) message(VERBOSE "Math test ${name} will be skipped as MPFR library is not available.") return() endif() diff --git a/libc/test/src/math/CMakeLists.txt b/libc/test/src/math/CMakeLists.txt index fcb47449748dc..87869538653f5 100644 --- a/libc/test/src/math/CMakeLists.txt +++ b/libc/test/src/math/CMakeLists.txt @@ -1722,3 +1722,8 @@ if(NOT LLVM_LIBC_FULL_BUILD) add_subdirectory(exhaustive) add_subdirectory(differential_testing) endif() + +# The GPU build uses special case exhaustive math tests. +if(LIBC_TARGET_ARCHITECTURE_IS_GPU) + add_subdirectory(gpu) +endif() diff --git a/libc/test/src/math/gpu/CMakeLists.txt b/libc/test/src/math/gpu/CMakeLists.txt new file mode 100644 index 0000000000000..43ddcfa452576 --- /dev/null +++ b/libc/test/src/math/gpu/CMakeLists.txt @@ -0,0 +1,134 @@ +add_libc_exhaustive_testsuite(libc-math-gpu-exhaustive-tests) + +if((NOT "openmp" IN_LIST LLVM_ENABLE_RUNTIMES) AND + (NOT "openmp" IN_LIST LLVM_ENABLE_PROJECTS)) + message(STATUS "The 'openmp' runtime must be enabled to run exhaustive " + "GPU tests.") + return() +endif() + +# Attempt to locate the libraries required for offloading. +if(TARGET omptarget.devicertl AND TARGET omptarget AND TARGET omp) + set(LIBC_OPENMP_RUNTIME omptarget.devicertl omptarget omp) +else() + find_library(omptarget.devicertl NAMES omptarget.devicertl + PATHS ${LLVM_LIBRARY_OUTPUT_INTDIR} ${LLVM_LIBRARY_DIR} + NO_DEFAULT_PATH + ) + find_library(omptarget NAMES omptarget + PATHS ${LLVM_LIBRARY_OUTPUT_INTDIR} ${LLVM_LIBRARY_DIR} + NO_DEFAULT_PATH + ) + find_library(omp NAMES omp + PATHS ${LLVM_LIBRARY_OUTPUT_INTDIR} ${LLVM_LIBRARY_DIR} + NO_DEFAULT_PATH + ) + if(NOT omptarget.devicertl OR NOT omptarget OR NOT omp) + message(WARNING "Could not find the OpenMP runtime for exhaustive tests") + return() + endif() + set(LIBC_OPENMP_RUNTIME ${omptarget.devicertl} ${omptarget} ${omp}) +endif() + +# Ensure that the tests do not use any other libraries found on the system. +if(${CMAKE_HOST_SYSTEM_NAME} MATCHES "Linux") + list(APPEND LIBC_OPENMP_RUNTIME "-Wl,-rpath,${LLVM_LIBRARY_DIR}" + "-Wl,-rpath,${LLVM_LIBRARY_OUTPUT_INTDIR}") +endif() + +function(add_gpu_exhaustive_test name) + cmake_parse_arguments( + "MATH_GPU_EXHAUSTIVE_TEST" + "NEED_MPFR" # Optional arguments + "" # Single value arguments + "SRCS;HDRS;DEPENDS;ARGS;ENV;COMPILE_OPTIONS;LINK_LIBRARIES" # Multi-value arguments + ${ARGN} + ) + + if(NOT MATH_GPU_EXHAUSTIVE_TEST_SRCS) + message(FATAL_ERROR "'add_gpu_exhaustive_test' target requires a SRCS list " + "of .cpp files.") + endif() + if(NOT MATH_GPU_EXHAUSTIVE_TEST_DEPENDS) + message(FATAL_ERROR "'add_gpu_exhaustive_test' target requires a DEPENDS " + "list of 'add_entrypoint_object' targets.") + endif() + + if(MATH_GPU_EXHAUSTIVE_TEST_NEED_MPFR) + if(NOT LIBC_TESTS_CAN_USE_MPFR) + message(VERBOSE "Math test ${name} will be skipped as MPFR library is " + "not available.") + return() + endif() + endif() + + + if(MATH_GPU_EXHAUSTIVE_TEST_NEED_MPFR) + list(APPEND MATH_GPU_EXHAUSTIVE_TEST_LINK_LIBRARIES libcMPFRWrapper) + endif() + + get_fq_target_name(${name} fq_target_name) + get_fq_deps_list(fq_deps_list ${MATH_GPU_EXHAUSTIVE_TEST_DEPENDS}) + list(REMOVE_DUPLICATES fq_deps_list) + + get_object_files_for_test( + link_object_files skipped_entrypoints_list ${fq_deps_list}) + if(skipped_entrypoints_list) + message(STATUS "Skipping unittest ${fq_target_name} as it has missing deps:" + " ${skipped_entrypoints_list}.") + return() + endif() + + list(APPEND MATH_GPU_EXHAUSTIVE_TEST_LINK_LIBRARIES + ${LIBC_OPENMP_RUNTIME} "--offload-link") + foreach(link_object_file ${link_object_files}) + list(APPEND MATH_GPU_EXHAUSTIVE_TEST_LINK_LIBRARIES + "-Xoffload-linker ${link_object_file}") + endforeach() + + set(fq_exhaustive_target_name ${fq_target_name}.__exhaustive__) + add_executable( + ${fq_exhaustive_target_name} + EXCLUDE_FROM_ALL + ${MATH_GPU_EXHAUSTIVE_TEST_SRCS} + ${MATH_GPU_EXHAUSTIVE_TEST_HDRS} + ) + set(LIBC_GPU_TEST_OPTIONS -fopenmp -nogpulib -nogpuinc -foffload-lto -fno-rtti + -fopenmp-offload-mandatory -fpie -fno-exceptions + --offload-arch=${LIBC_GPU_TARGET_ARCHITECTURE}) + + target_include_directories(${fq_exhaustive_target_name} PRIVATE + ${LIBC_SOURCE_DIR}) + + target_compile_options(${fq_exhaustive_target_name} PRIVATE + ${MATH_GPU_EXHAUSTIVE_TEST_COMPILE_OPTIONS} + ${LIBC_GPU_TEST_OPTIONS} + ) + target_link_libraries(${fq_exhaustive_target_name} PRIVATE + ${MATH_GPU_EXHAUSTIVE_TEST_LINK_LIBRARIES} LibcTest.unit + ) + + add_dependencies(${fq_exhaustive_target_name} ${fq_deps_list}) + set_target_properties(${fq_exhaustive_target_name} + PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) + + add_custom_target( + ${fq_target_name} + COMMAND OMP_TARGET_OFFLOAD=MANDATORY + $ + COMMAND_EXPAND_LISTS + COMMENT "Running exhaustive GPU test ${fq_target_name}" + ) + add_dependencies(libc-math-gpu-exhaustive-tests ${fq_target_name}) +endfunction() + +add_gpu_exhaustive_test( + truncf_test + NEED_MPFR + SRCS + truncf_test.cpp + DEPENDS + libc.include.math + libc.src.math.truncf + libc.src.__support.FPUtil.fp_bits +) diff --git a/libc/test/src/math/gpu/exhaustive_test.h b/libc/test/src/math/gpu/exhaustive_test.h new file mode 100644 index 0000000000000..73965eac3f070 --- /dev/null +++ b/libc/test/src/math/gpu/exhaustive_test.h @@ -0,0 +1,137 @@ +//===-- Exhaustive test template for math functions -------------*- C++ -*-===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include "src/__support/CPP/type_traits.h" +#include "src/__support/FPUtil/FPBits.h" +#include "test/UnitTest/FPMatcher.h" +#include "test/UnitTest/Test.h" +#include "utils/MPFRWrapper/MPFRUtils.h" + +#include +#include +#include +#include +#include +#include +#include + +// To test exhaustively for inputs in the range [start, stop) in parallel: +// 1. Define a Checker class with: +// - FloatType: define floating point type to be used. +// - FPBits: fputil::FPBits. +// - UIntType: define bit type for the corresponding floating point type. +// - uint64_t check(start, stop, rounding_mode): a method to test in given +// range for a given rounding mode, which returns the number of +// failures. +// 2. Use LlvmLibcExhaustiveMathTest class +// 3. Call: test_full_range(start, stop, nthreads, rounding) +// or test_full_range_all_roundings(start, stop). +// * For single input single output math function, use the convenient template: +// LlvmLibcUnaryOpExhaustiveMathTest. +namespace mpfr = LIBC_NAMESPACE::testing::mpfr; + +template using UnaryOp = T(T); + +template Func> +struct UnaryOpChecker : public virtual LIBC_NAMESPACE::testing::Test { + using FloatType = T; + using FPBits = LIBC_NAMESPACE::fputil::FPBits; + using UIntType = typename FPBits::UIntType; + + static constexpr UnaryOp *FUNC = Func; + static constexpr mpfr::Operation OP = Op; + + // Check in a range, return the number of failures. + bool check(FloatType in, FloatType out, mpfr::RoundingMode rounding) { + mpfr::ForceRoundingMode r(rounding); + if (!r.success) + return true; + + bool correct = TEST_MPFR_MATCH_ROUNDING(Op, in, out, 0.5, rounding); + return !correct; + } +}; + +// Checker class needs inherit from LIBC_NAMESPACE::testing::Test and provide +// UIntType and check method. +template +struct LlvmLibcExhaustiveMathTest + : public virtual LIBC_NAMESPACE::testing::Test, + public Checker { + using FloatType = typename Checker::FloatType; + using FPBits = typename Checker::FPBits; + using UIntType = typename Checker::UIntType; + + static constexpr UIntType BLOCK_SIZE = (1 << 25); + + // Break [start, stop) into chunks and compare results on the GPU vs the CPU. + void test_full_range(UIntType start, UIntType stop, + mpfr::RoundingMode rounding) { + + // TODO: We can run the GPU asynchronously to compute the next block. + // However, the main bottleneck is MPFR on the CPU. + uint64_t failed = 0; + for (UIntType chunk = start; chunk <= stop; chunk += BLOCK_SIZE) { + uint64_t percent = (static_cast(chunk - start) / + static_cast(stop - start)) * + 100.0; + std::cout << percent << "% is in process \r" << std::flush; + UIntType end = std::min(stop, chunk + BLOCK_SIZE); + + std::vector data(BLOCK_SIZE, FloatType(0)); + + FloatType *ptr = data.data(); + // Fill the buffer with the computed results from the GPU. +#pragma omp target teams distribute parallel for map(from : ptr[0 : BLOCK_SIZE]) + for (UIntType begin = chunk; begin < end; ++begin) { + UIntType idx = begin - chunk; + + FPBits xbits(begin); + FloatType x = FloatType(xbits); + + ptr[idx] = Checker::FUNC(x); + } + + std::atomic failed_in_range = 0; + // Check the GPU results against the MPFR library. +#pragma omp parallel for default(firstprivate) shared(failed_in_range) + for (UIntType begin = chunk; begin < end; ++begin) { + UIntType idx = begin - chunk; + + FPBits xbits(begin); + FloatType x = FloatType(xbits); + + failed_in_range += Checker::check(x, data[idx], rounding); + } + + if (failed_in_range > 0) { + std::stringstream msg; + msg << "Test failed for " << std::dec << failed_in_range + << " inputs in range: " << chunk << " to " << end << " [0x" + << std::hex << chunk << ", 0x" << end << "), [" << std::hexfloat + << static_cast(FPBits(chunk)) << ", " + << static_cast(FPBits(end)) << ")\n"; + std::cerr << msg.str() << std::flush; + + failed += failed_in_range.load(); + } + + // Check to make sure we don't overflow when updating the value. + if (chunk > std::numeric_limits::max() - BLOCK_SIZE) + chunk = std::numeric_limits::max(); + } + + std::cout << std::endl; + std::cout << "Test " << ((failed > 0) ? "FAILED" : "PASSED") << std::endl; + ASSERT_EQ(failed, uint64_t(0)); + } +}; + +template Func> +using LlvmLibcUnaryOpExhaustiveMathTest = + LlvmLibcExhaustiveMathTest>; diff --git a/libc/test/src/math/gpu/truncf_test.cpp b/libc/test/src/math/gpu/truncf_test.cpp new file mode 100644 index 0000000000000..c708c544522f0 --- /dev/null +++ b/libc/test/src/math/gpu/truncf_test.cpp @@ -0,0 +1,33 @@ +//===-- Exhaustive GPU test for truncf ------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include "exhaustive_test.h" +#include "src/math/truncf.h" +#include "utils/MPFRWrapper/MPFRUtils.h" + +namespace mpfr = LIBC_NAMESPACE::testing::mpfr; + +using LlvmLibcTruncfExhaustiveTest = + LlvmLibcUnaryOpExhaustiveMathTest; + +// Range: [0, Inf]; +static constexpr uint32_t POS_START = 0x0000'0000U; +static constexpr uint32_t POS_STOP = 0x7f80'0000U; + +TEST_F(LlvmLibcTruncfExhaustiveTest, PostiveRange) { + test_full_range(POS_START, POS_STOP, mpfr::RoundingMode::Nearest); +} + +// Range: [-Inf, 0]; +static constexpr uint32_t NEG_START = 0xb000'0000U; +static constexpr uint32_t NEG_STOP = 0xff80'0000U; + +TEST_F(LlvmLibcTruncfExhaustiveTest, NegativeRange) { + test_full_range(NEG_START, NEG_STOP, mpfr::RoundingMode::Nearest); +} diff --git a/libc/utils/MPFRWrapper/CMakeLists.txt b/libc/utils/MPFRWrapper/CMakeLists.txt index 416307a20d7d1..7d3b46df369ca 100644 --- a/libc/utils/MPFRWrapper/CMakeLists.txt +++ b/libc/utils/MPFRWrapper/CMakeLists.txt @@ -21,6 +21,6 @@ if(LIBC_TESTS_CAN_USE_MPFR) target_link_directories(libcMPFRWrapper PUBLIC ${LLVM_LIBC_MPFR_INSTALL_PATH}/lib) endif() target_link_libraries(libcMPFRWrapper LibcFPTestHelpers.unit LibcTest.unit mpfr gmp) -elseif(NOT LIBC_TARGET_ARCHITECTURE_IS_GPU) +else() message(WARNING "Math tests using MPFR will be skipped.") endif()