Fix cmake folder
ADD
cmake/VerifyCompiler.cmake
+# MIT License
+#
+# Copyright (c) 2019-2020 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.
+
+list(APPEND CMAKE_PREFIX_PATH /opt/rocm /opt/rocm/hip)
+find_package(hip REQUIRED CONFIG PATHS /opt/rocm)
+
+if(HIP_COMPILER STREQUAL "nvcc")
+ message(FATAL_ERROR "rocThrust does not support the CUDA backend.")
+elseif(HIP_COMPILER STREQUAL "hcc" OR HIP_COMPILER STREQUAL "clang")
+ if(NOT (CMAKE_CXX_COMPILER MATCHES ".*/hcc$" OR CMAKE_CXX_COMPILER MATCHES ".*/hipcc$"))
+ message(FATAL_ERROR "On ROCm platform 'hcc' or 'clang' must be used as C++ compiler.")
+ elseif(NOT CXX_VERSION_STRING MATCHES "clang")
+ list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hcc)
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unused-command-line-argument")
+ endif()
+
+ if(HIP_COMPILER STREQUAL "hcc")
+ list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hcc)
+ find_package(hcc REQUIRED CONFIG PATHS /opt/rocm)
+ endif()
+else()
+ message(FATAL_ERROR "HIP_COMPILER must be 'hcc' or 'clang' (AMD ROCm platform)")
+endif()

DELETE
cmake/header_test.in
-// This source file checks that:
-// 1) Header <thrust/${header}> compiles without error.
-// 2) Common macro collisions with platform/system headers are avoided.
-
-// Turn off failures for certain configurations:
-#define THRUST_CPP11_REQUIRED_NO_ERROR
-#define THRUST_CPP14_REQUIRED_NO_ERROR
-#define THRUST_MODERN_GCC_REQUIRED_NO_ERROR
-
-#ifndef THRUST_IGNORE_MACRO_CHECKS
-
-// Define THRUST_MACRO_CHECK(macro, header), which emits a diagnostic indicating
-// a potential macro collision and halts.
-//
-// Hacky way to build a string, but it works on all tested platforms.
-#define THRUST_MACRO_CHECK(MACRO, HEADER) \
- THRUST_MACRO_CHECK_IMPL(Identifier MACRO should not be used from Thrust \
- headers due to conflicts with HEADER.)
-
-// Use raw platform checks instead of the THRUST_HOST_COMPILER macros since we
-// don't want to #include any headers other than the one being tested.
-//
-// This is only implemented for MSVC/GCC/Clang.
-#if defined(_MSC_VER) // MSVC
-
-// Fake up an error for MSVC
-#define THRUST_MACRO_CHECK_IMPL(msg) \
- /* Print message that looks like an error: */ \
- __pragma(message(__FILE__ ":" THRUST_MACRO_CHECK_IMPL0(__LINE__) \
- ": error: " #msg)) \
- /* abort compilation due to static_assert or syntax error: */ \
- static_assert(false, #msg);
-#define THRUST_MACRO_CHECK_IMPL0(x) THRUST_MACRO_CHECK_IMPL1(x)
-#define THRUST_MACRO_CHECK_IMPL1(x) #x
-
-#elif defined(__clang__) || defined(__GNUC__)
-
-// GCC/clang are easy:
-#define THRUST_MACRO_CHECK_IMPL(msg) THRUST_MACRO_CHECK_IMPL0(GCC error #msg)
-#define THRUST_MACRO_CHECK_IMPL0(expr) _Pragma(#expr)
-
-#endif
-
-// complex.h conflicts
-#define I THRUST_MACRO_CHECK('I', complex.h)
-
-// windows.h conflicts
-// Disabling for now; we use min/max in many places, but since most
-// projects build with NOMINMAX this doesn't seem to be high priority to fix.
-//#define min(...) THRUST_MACRO_CHECK('min', windows.h)
-//#define max(...) THRUST_MACRO_CHECK('max', windows.h)
-
-#endif // THRUST_IGNORE_MACRO_CHECKS
-
-#include <thrust/${header}>

MODIFY
test/test_scan.cpp
test/test_scan.cpp
@@ -312,7 +312,7 @@ void TestScanMixedTypes(void)
thrust::inclusive_scan(float_input.begin(), float_input.end(), int_output.begin());
ASSERT_EQ(int_output[0], 1); // in: 1.5 accum: 1.5f out: 1
ASSERT_EQ(int_output[1], 4); // in: 2.5 accum: 4.0f out: 4
- ASSERT_EQL(int_output[2], 7); // in: 3.5 accum: 7.5f out: 7
+ ASSERT_EQ(int_output[2], 7); // in: 3.5 accum: 7.5f out: 7
ASSERT_EQ(int_output[3], 12); // in: 4.5 accum: 12.f out: 12
// float -> float with plus<int> operator (float accumulator)

MODIFY
thrust/system/hip/detail/cross_system.h
thrust/system/hip/detail/cross_system.h
@@ -141,6 +141,7 @@ struct cross_system : execution_policy<cross_system<Sys1, Sys2> >
// MSVC2015 WAR: put decltype here instead of in trailing return type
typename Direction =
decltype(direction_of_copy(std::declval<ExecutionPolicy>()))>
+ THRUST_CONSTEXPR __host__ __device__
auto is_device_to_host_copy(ExecutionPolicy const& exec)
noexcept ->
thrust::detail::integral_constant<

MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -52,9 +52,13 @@ namespace thrust
namespace hip_rocprim
{
- __host__ __device__ inline hipStream_t default_stream()
+ inline __host__ __device__ hipStream_t default_stream()
{
+ #ifdef HIP_API_PER_THREAD_DEFAULT_STREAM
+ return hipStreamPerThread;
+ #else
return hipStreamDefault; // There's not hipStreamLegacy
+ #endif
}
template <class Derived>
@@ -179,6 +183,12 @@ inline void __host__ __device__ terminate()
inline void __host__ __device__ throw_on_error(hipError_t status, char const* msg)
{
+#if __THRUST_HAS_HIPRT__
+ // Clear the global HIP error state which may have been set by the last
+ // call. Otherwise, errors may "leak" to unrelated kernel launches.
+ hipGetLastError();
+ #endif
+
if(hipSuccess != status)
{
if (THRUST_IS_HOST_CODE) {
@@ -193,6 +203,7 @@ inline void __host__ __device__ throw_on_error(hipError_t status, char const* ms
THRUST_HIP_PRINTF("Error %d :%s \n", (int)status, msg);
#if THRUST_HIP_PRINTF_ENABLED == 0
THRUST_UNUSED_VAR(status);
+ THRUST_UNUSED_VAR(msg);
#endif
#endif
hip_rocprim::terminate();
@@ -204,6 +215,12 @@ inline void __host__ __device__ throw_on_error(hipError_t status, char const* ms
// TODO this overload should be removed and messages should be passed.
inline void __host__ __device__ throw_on_error(hipError_t status)
{
+#if __THRUST_HAS_HIPRT__
+ // Clear the global HIP error state which may have been set by the last
+ // call. Otherwise, errors may "leak" to unrelated kernel launches.
+ hipGetLastError();
+ #endif
+
if(hipSuccess != status)
{
if (THRUST_IS_HOST_CODE) {

Fix cmake folder