From dd2d0e7c0f9353399c57b1f3906227200e9eb1c9 Mon Sep 17 00:00:00 2001 From: Saad Rahim Date: Wed, 8 May 2019 13:48:40 -0700 Subject: [PATCH 1/4] Adding NOTICES.txt --- NOTICES.txt | 67 +++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 67 insertions(+) create mode 100644 NOTICES.txt diff --git a/NOTICES.txt b/NOTICES.txt new file mode 100644 index 000000000..dccc68ab2 --- /dev/null +++ b/NOTICES.txt @@ -0,0 +1,67 @@ +Notices and Licenses file +______________________________________________________________________________ + +AMD copyrighted code (MIT) +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. + + +ROCmSoftwarePlatform-rocPRIM v2.5.0 (MIT) +Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. + + +florianrappl-cmdparser v-u (MIT) +Copyright (c) 2015 - 2016 Florian Rappl + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. From 9c288c2b43a3e5f471a88c171fe9418c9ecb7a78 Mon Sep 17 00:00:00 2001 From: Saad Rahim Date: Wed, 8 May 2019 16:54:21 -0700 Subject: [PATCH 2/4] Fixing dockerfile --- docker/dockerfile-build-ubuntu-rock | 11 ----------- docker/dockerfile-install-ubuntu | 10 ---------- 2 files changed, 21 deletions(-) diff --git a/docker/dockerfile-build-ubuntu-rock b/docker/dockerfile-build-ubuntu-rock index fc06aa089..e8a463361 100644 --- a/docker/dockerfile-build-ubuntu-rock +++ b/docker/dockerfile-build-ubuntu-rock @@ -37,14 +37,3 @@ RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-ins RUN useradd --create-home -u ${user_uid} -o -G sudo --shell /bin/bash jenkins && \ mkdir -p /etc/sudoers.d/ && \ echo '%sudo ALL=(ALL) NOPASSWD:ALL' | tee /etc/sudoers.d/sudo-nopasswd - -ARG ROCBLAS_SRC_ROOT=/usr/local/src/rocBLAS - -# Clone rocblas repo -# Build client dependencies and install into /usr/local (LAPACK & GTEST) -RUN mkdir -p ${ROCBLAS_SRC_ROOT} && cd ${ROCBLAS_SRC_ROOT} && \ - git clone -b saad/develop --depth=1 https://github.com/saadrahim/rocPRIM . #&& \ - #mkdir -p build/deps && cd build/deps && \ - #cmake -DBUILD_BOOST=OFF ${ROCBLAS_SRC_ROOT}/deps && \ - #make -j $(nproc) install && \ - #rm -rf ${ROCBLAS_SRC_ROOT} diff --git a/docker/dockerfile-install-ubuntu b/docker/dockerfile-install-ubuntu index 7d20aad61..45e329901 100755 --- a/docker/dockerfile-install-ubuntu +++ b/docker/dockerfile-install-ubuntu @@ -42,13 +42,3 @@ RUN useradd --create-home -u ${user_uid} -o -G sudo --shell /bin/bash jenkins && mkdir -p /etc/sudoers.d/ && \ echo '%sudo ALL=(ALL) NOPASSWD:ALL' | tee /etc/sudoers.d/sudo-nopasswd -ARG ROCBLAS_SRC_ROOT=/usr/local/src/rocBLAS - -# Clone rocblas repo -# Build client dependencies and install into /usr/local (LAPACK & GTEST) -RUN mkdir -p ${ROCBLAS_SRC_ROOT} && cd ${ROCBLAS_SRC_ROOT} && \ - git clone -b develop --depth=1 https://github.com/ROCmSoftwarePlatform/rocBLAS . && \ - mkdir -p build/deps && cd build/deps && \ - cmake -DBUILD_BOOST=OFF ${ROCBLAS_SRC_ROOT}/deps && \ - make -j $(nproc) install -# rm -rf ${ROCBLAS_SRC_ROOT} From 351a2f7326cb5250bb912b929b7c243da2810432 Mon Sep 17 00:00:00 2001 From: Anton Gorenko Date: Tue, 14 May 2019 16:48:35 +0600 Subject: [PATCH 3/4] Implement invoke_result for device-only lambdas invoke_result is based on https://en.cppreference.com/w/cpp/types/result_of The main difference is using ROCPRIM_HOST_DEVICE, this allows to use invoke_result with device-only lambdas/functors in host-only functions on HIP-clang. Workarounds with ROCPRIM_HOST_DEVICE for such lambdas are removed. --- .../rocprim/detail/match_result_type.hpp | 85 ++++++++++++++----- .../device/detail/device_transform.hpp | 7 +- .../rocprim/device/device_binary_search.hpp | 5 -- .../rocprim/device/device_scan_by_key.hpp | 5 -- .../rocprim/device/device_segmented_scan.hpp | 5 -- .../rocprim/device/device_transform.hpp | 7 +- .../rocprim/iterator/transform_iterator.hpp | 10 +-- 7 files changed, 72 insertions(+), 52 deletions(-) diff --git a/rocprim/include/rocprim/detail/match_result_type.hpp b/rocprim/include/rocprim/detail/match_result_type.hpp index 75ba31b01..e143a8aac 100644 --- a/rocprim/include/rocprim/detail/match_result_type.hpp +++ b/rocprim/include/rocprim/detail/match_result_type.hpp @@ -24,38 +24,85 @@ #include #include "../config.hpp" -#include "../types/tuple.hpp" BEGIN_ROCPRIM_NAMESPACE namespace detail { -// tuple_contains_type::value is false if Tuple is not rocprim::tuple<> or Tuple is -// rocprim::tuple<> class which does not contain element of type T; otherwise it's true. -template -struct tuple_contains_type : std::false_type {}; +// invoke_result is based on https://en.cppreference.com/w/cpp/types/result_of +// The main difference is using ROCPRIM_HOST_DEVICE, this allows to +// use invoke_result with device-only lambdas/functors in host-only functions +// on HIP-clang. + +template +struct is_reference_wrapper : std::false_type {}; +template +struct is_reference_wrapper> : std::true_type {}; template -struct tuple_contains_type> : std::false_type {}; +struct invoke_impl { + template + ROCPRIM_HOST_DEVICE + static auto call(F&& f, Args&&... args) + -> decltype(std::forward(f)(std::forward(args)...)); +}; + +template +struct invoke_impl +{ + template::type, + class = typename std::enable_if::value>::type + > + ROCPRIM_HOST_DEVICE + static auto get(T&& t) -> T&&; + + template::type, + class = typename std::enable_if::value>::type + > + ROCPRIM_HOST_DEVICE + static auto get(T&& t) -> decltype(t.get()); + + template::type, + class = typename std::enable_if::value>::type, + class = typename std::enable_if::value>::type + > + ROCPRIM_HOST_DEVICE + static auto get(T&& t) -> decltype(*std::forward(t)); -template -struct tuple_contains_type> : tuple_contains_type> {}; + template::value>::type + > + ROCPRIM_HOST_DEVICE + static auto call(MT1 B::*pmf, T&& t, Args&&... args) + -> decltype((invoke_impl::get(std::forward(t)).*pmf)(std::forward(args)...)); + + template + ROCPRIM_HOST_DEVICE + static auto call(MT B::*pmd, T&& t) + -> decltype(invoke_impl::get(std::forward(t)).*pmd); +}; + +template::type> +ROCPRIM_HOST_DEVICE +auto INVOKE(F&& f, Args&&... args) + -> decltype(invoke_impl::call(std::forward(f), std::forward(args)...)); + +// Conforming C++14 implementation (is also a valid C++11 implementation): +template +struct invoke_result_impl { }; +template +struct invoke_result_impl(), std::declval()...))), F, Args...> +{ + using type = decltype(INVOKE(std::declval(), std::declval()...)); +}; -template -struct tuple_contains_type> : std::true_type {}; +template +struct invoke_result : invoke_result_impl {}; template struct match_result_type { -private: - #ifdef __cpp_lib_is_invocable - using binary_result_type = typename std::invoke_result::type; - #else - using binary_result_type = typename std::result_of::type; - #endif - -public: - using type = binary_result_type; + using type = typename invoke_result::type; }; } // end namespace detail diff --git a/rocprim/include/rocprim/device/detail/device_transform.hpp b/rocprim/include/rocprim/device/detail/device_transform.hpp index 24febc211..3665daded 100644 --- a/rocprim/include/rocprim/device/detail/device_transform.hpp +++ b/rocprim/include/rocprim/device/detail/device_transform.hpp @@ -26,6 +26,7 @@ #include "../../config.hpp" #include "../../detail/various.hpp" +#include "../../detail/match_result_type.hpp" #include "../../intrinsics.hpp" #include "../../functional.hpp" @@ -44,11 +45,7 @@ namespace detail template struct unpack_binary_op { - #ifdef __cpp_lib_is_invocable - using result_type = typename std::invoke_result::type; - #else - using result_type = typename std::result_of::type; - #endif + using result_type = typename ::rocprim::detail::invoke_result::type; ROCPRIM_HOST_DEVICE inline unpack_binary_op() = default; diff --git a/rocprim/include/rocprim/device/device_binary_search.hpp b/rocprim/include/rocprim/device/device_binary_search.hpp index 8cdfd2a64..5f655b227 100644 --- a/rocprim/include/rocprim/device/device_binary_search.hpp +++ b/rocprim/include/rocprim/device/device_binary_search.hpp @@ -74,12 +74,7 @@ hipError_t binary_search(void * temporary_storage, needles, output, needles_size, [haystack, haystack_size, search_op, compare_op] - #ifdef __HIP__ - // Workaround: hip-clang does not support std::result_of of device-only functions - ROCPRIM_HOST_DEVICE - #else ROCPRIM_DEVICE - #endif (const value_type& value) { return search_op(haystack, haystack_size, value, compare_op); diff --git a/rocprim/include/rocprim/device/device_scan_by_key.hpp b/rocprim/include/rocprim/device/device_scan_by_key.hpp index 97dad64e6..b9cee7231 100644 --- a/rocprim/include/rocprim/device/device_scan_by_key.hpp +++ b/rocprim/include/rocprim/device/device_scan_by_key.hpp @@ -164,12 +164,7 @@ hipError_t inclusive_scan_by_key(void * temporary_storage, rocprim::make_transform_iterator( rocprim::make_counting_iterator(0), [values_input, keys_input, key_compare_op] - #ifdef __HIP__ - // Workaround: hip-clang does not support std::result_of of device-only functions - ROCPRIM_HOST_DEVICE - #else ROCPRIM_DEVICE - #endif (const size_t i) { flag_type flag(true); diff --git a/rocprim/include/rocprim/device/device_segmented_scan.hpp b/rocprim/include/rocprim/device/device_segmented_scan.hpp index e171c518b..7c3a9e967 100644 --- a/rocprim/include/rocprim/device/device_segmented_scan.hpp +++ b/rocprim/include/rocprim/device/device_segmented_scan.hpp @@ -614,12 +614,7 @@ hipError_t segmented_exclusive_scan(void * temporary_storage, rocprim::make_transform_iterator( rocprim::make_counting_iterator(0), [input, head_flags, initial_value_converted, size] - #ifdef __HIP__ - // Workaround: hip-clang does not support std::result_of of device-only functions - ROCPRIM_HOST_DEVICE - #else ROCPRIM_DEVICE - #endif (const size_t i) { flag_type flag(false); diff --git a/rocprim/include/rocprim/device/device_transform.hpp b/rocprim/include/rocprim/device/device_transform.hpp index 2779a85a6..bccf3fb78 100644 --- a/rocprim/include/rocprim/device/device_transform.hpp +++ b/rocprim/include/rocprim/device/device_transform.hpp @@ -26,6 +26,7 @@ #include "../config.hpp" #include "../detail/various.hpp" +#include "../detail/match_result_type.hpp" #include "../types/tuple.hpp" #include "../iterator/zip_iterator.hpp" @@ -145,11 +146,7 @@ hipError_t transform(InputIterator input, bool debug_synchronous = false) { using input_type = typename std::iterator_traits::value_type; - #ifdef __cpp_lib_is_invocable - using result_type = typename std::invoke_result::type; - #else - using result_type = typename std::result_of::type; - #endif + using result_type = typename ::rocprim::detail::invoke_result::type; // Get default config if Config is default_config using config = detail::default_or_custom_config< diff --git a/rocprim/include/rocprim/iterator/transform_iterator.hpp b/rocprim/include/rocprim/iterator/transform_iterator.hpp index d53c20738..2880a1ea0 100644 --- a/rocprim/include/rocprim/iterator/transform_iterator.hpp +++ b/rocprim/include/rocprim/iterator/transform_iterator.hpp @@ -26,6 +26,7 @@ #include #include "../config.hpp" +#include "../detail/match_result_type.hpp" /// \addtogroup iteratormodule /// @{ @@ -49,17 +50,10 @@ BEGIN_ROCPRIM_NAMESPACE template< class InputIterator, class UnaryFunction, -#if defined(__cpp_lib_is_invocable) && !defined(DOXYGEN_SHOULD_SKIP_THIS) // C++17 class ValueType = - typename std::invoke_result< + typename ::rocprim::detail::invoke_result< UnaryFunction, typename std::iterator_traits::value_type >::type -#else - class ValueType = - typename std::result_of< - UnaryFunction(typename std::iterator_traits::value_type) - >::type -#endif > class transform_iterator { From 9962bf8285fd4b1e1c624c0b4d29782b465cc9ae Mon Sep 17 00:00:00 2001 From: saadrahim <44449863+saadrahim@users.noreply.github.com> Date: Wed, 15 May 2019 14:26:37 -0600 Subject: [PATCH 4/4] Version number based on rocm release (#69) * Version number based on rocm release * Fixing whitespace * Fixing based on review for PR --- CMakeLists.txt | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3969a4c69..6acc581c4 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,7 +26,7 @@ cmake_minimum_required(VERSION 3.5.1 FATAL_ERROR) set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "Install path prefix, prepended onto install directories") # rocPRIM project -project(rocprim VERSION 1.0.2.0 LANGUAGES CXX) +project(rocprim LANGUAGES CXX) # CMake modules list(APPEND CMAKE_MODULE_PATH @@ -72,6 +72,17 @@ include(cmake/Dependencies.cmake) # AMD targets set(AMDGPU_TARGETS gfx803;gfx900;gfx906 CACHE STRING "List of specific machine types for library to target") +# Setup VERSION +set ( VERSION_STRING "2.5.0." ) + +# Check if BUILD_NUMBER is defined in a Jenkins envirnment +if ($ENV{BUILD_NUMBER} ) + string(CONCAT BUILD_VERSION ${VERSION_STRING} $ENV{BUILD_NUMBER}) +else () + string(CONCAT BUILD_VERSION ${VERSION_STRING} "0") +endif () +rocm_setup_version( VERSION ${BUILD_VERSION} NO_GIT_TAG_VERSION ) + # Print configuration summary include(cmake/Summary.cmake) print_configuration_summary()