Merge pull request #71 from ROCmSoftwarePlatform/develop

Develop to Master

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()

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.

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}

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}

rocprim/include/rocprim/detail/match_result_type.hpp

@@ -24,38 +24,85 @@
 #include <type_traits>
 
 #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<class T, class Tuple>
-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 <class T>
+struct is_reference_wrapper : std::false_type {};
+template <class U>
+struct is_reference_wrapper<std::reference_wrapper<U>> : std::true_type {};
 
 template<class T>
-struct tuple_contains_type<T, ::rocprim::tuple<>> : std::false_type {};
+struct invoke_impl {
+    template<class F, class... Args>
+    ROCPRIM_HOST_DEVICE
+    static auto call(F&& f, Args&&... args)
+        -> decltype(std::forward<F>(f)(std::forward<Args>(args)...));
+};
+
+template<class B, class MT>
+struct invoke_impl<MT B::*>
+{
+    template<class T, class Td = typename std::decay<T>::type,
+        class = typename std::enable_if<std::is_base_of<B, Td>::value>::type
+    >
+    ROCPRIM_HOST_DEVICE
+    static auto get(T&& t) -> T&&;
+
+    template<class T, class Td = typename std::decay<T>::type,
+        class = typename std::enable_if<is_reference_wrapper<Td>::value>::type
+    >
+    ROCPRIM_HOST_DEVICE
+    static auto get(T&& t) -> decltype(t.get());
+
+    template<class T, class Td = typename std::decay<T>::type,
+        class = typename std::enable_if<!std::is_base_of<B, Td>::value>::type,
+        class = typename std::enable_if<!is_reference_wrapper<Td>::value>::type
+    >
+    ROCPRIM_HOST_DEVICE
+    static auto get(T&& t) -> decltype(*std::forward<T>(t));
 
-template<class T, class U, class... Ts>
-struct tuple_contains_type<T, ::rocprim::tuple<U, Ts...>> : tuple_contains_type<T, ::rocprim::tuple<Ts...>> {};
+    template<class T, class... Args, class MT1,
+        class = typename std::enable_if<std::is_function<MT1>::value>::type
+    >
+    ROCPRIM_HOST_DEVICE
+    static auto call(MT1 B::*pmf, T&& t, Args&&... args)
+        -> decltype((invoke_impl::get(std::forward<T>(t)).*pmf)(std::forward<Args>(args)...));
+
+    template<class T>
+    ROCPRIM_HOST_DEVICE
+    static auto call(MT B::*pmd, T&& t)
+        -> decltype(invoke_impl::get(std::forward<T>(t)).*pmd);
+};
+
+template<class F, class... Args, class Fd = typename std::decay<F>::type>
+ROCPRIM_HOST_DEVICE
+auto INVOKE(F&& f, Args&&... args)
+    -> decltype(invoke_impl<Fd>::call(std::forward<F>(f), std::forward<Args>(args)...));
+
+// Conforming C++14 implementation (is also a valid C++11 implementation):
+template <typename AlwaysVoid, typename, typename...>
+struct invoke_result_impl { };
+template <typename F, typename...Args>
+struct invoke_result_impl<decltype(void(INVOKE(std::declval<F>(), std::declval<Args>()...))), F, Args...>
+{
+    using type = decltype(INVOKE(std::declval<F>(), std::declval<Args>()...));
+};
 
-template<class T, class... Ts>
-struct tuple_contains_type<T, ::rocprim::tuple<T, Ts...>> : std::true_type {};
+template <class F, class... ArgTypes>
+struct invoke_result : invoke_result_impl<void, F, ArgTypes...> {};
 
 template<class InputType, class BinaryFunction>
 struct match_result_type
 {
-private:
-    #ifdef __cpp_lib_is_invocable
-    using binary_result_type = typename std::invoke_result<BinaryFunction, InputType, InputType>::type;
-    #else
-    using binary_result_type = typename std::result_of<BinaryFunction(InputType, InputType)>::type;
-    #endif
-
-public:
-    using type = binary_result_type;
+    using type = typename invoke_result<BinaryFunction, InputType, InputType>::type;
 };
 
 } // end namespace detail

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<class T1, class T2, class BinaryFunction>
 struct unpack_binary_op
 {
-    #ifdef __cpp_lib_is_invocable
-    using result_type = typename std::invoke_result<BinaryFunction, T1, T2>::type;
-    #else
-    using result_type = typename std::result_of<BinaryFunction(T1, T2)>::type;
-    #endif
+    using result_type = typename ::rocprim::detail::invoke_result<BinaryFunction, T1, T2>::type;
 
     ROCPRIM_HOST_DEVICE inline
     unpack_binary_op() = default;

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);

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<size_t>(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);

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<size_t>(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);

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<InputIterator>::value_type;
-    #ifdef __cpp_lib_is_invocable
-    using result_type = typename std::invoke_result<UnaryFunction, input_type>::type;
-    #else
-    using result_type = typename std::result_of<UnaryFunction(input_type)>::type;
-    #endif
+    using result_type = typename ::rocprim::detail::invoke_result<UnaryFunction, input_type>::type;
 
     // Get default config if Config is default_config
     using config = detail::default_or_custom_config<

rocprim/include/rocprim/iterator/transform_iterator.hpp

@@ -26,6 +26,7 @@
 #include <type_traits>
 
 #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<InputIterator>::value_type
         >::type
-#else
-    class ValueType =
-        typename std::result_of<
-            UnaryFunction(typename std::iterator_traits<InputIterator>::value_type)
-        >::type
-#endif
 >
 class transform_iterator
 {