Update system hip
MODIFY
CMakeLists.txt
CMakeLists.txt
@@ -94,10 +94,6 @@ if ("TBB" STREQUAL "${THRUST_DEVICE_SYSTEM}")
set (CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ${TBB_LD_FLAGS}")
set (THRUST_ADDITIONAL_LIBRARIES "${TBB_LIBRARIES}")
endif ()
-
- # There's a ton of these in the TBB backend, even though the code is correct.
- # TODO: silence these warnings in code instead
- append_option_if_available("-Wno-unused-parameter" THRUST_CXX_WARNINGS)
endif ()
if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
@@ -256,4 +252,3 @@ foreach (THRUST_TEST_SOURCE IN LISTS THRUST_TESTS)
add_test(run.${THRUST_TEST} ${THRUST_TEST})
add_test(run.${THRUST_TEST_RDC} ${THRUST_TEST_RDC})
endforeach ()
-

DELETE
internal/test/thrust.example.cuda.global_device_vector.filecheck
internal/test/thrust.example.cuda.global_device_vector.filecheck
DELETE
internal/test/thrust.example.scan_matrix_by_rows.filecheck
internal/test/thrust.example.scan_matrix_by_rows.filecheck
DELETE
internal/test/warningstester.cu
-//#include "cuda_runtime_api.h"
-#include "warningstester.h"
-
-int main()
-{
- return 0;
-}
-

MODIFY
thrust/system/hip/config.h
thrust/system/hip/config.h
#include <thrust/detail/config.h>
-#define THRUST_UNUSED_VAR(expr) do { (void)(expr); } while (0)
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP
#ifndef __HIP_DEVICE_COMPILE__
#define __THRUST_HAS_HIPRT__ 1

ADD
thrust/system/hip/detail/dispatch.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ * Modifications Copyright (c) 2020, Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/preprocessor.h>
+
+/**
+ * Dispatch between 32-bit and 64-bit index based versions of the same algorithm
+ * implementation. This version assumes that callables for both branches consist
+ * of the same tokens, and is intended to be used with Thrust-style dispatch
+ * interfaces, that always deduce the size type from the arguments.
+ */
+#define THRUST_INDEX_TYPE_DISPATCH(status, call, count, arguments) \
+ if (count <= std::numeric_limits<thrust::detail::int32_t>::max()) { \
+ thrust::detail::int32_t THRUST_PP_CAT2(count, _fixed) = count; \
+ status = call arguments; \
+ } \
+ else { \
+ thrust::detail::int64_t THRUST_PP_CAT2(count, _fixed) = count; \
+ status = call arguments; \
+ }
+
+/**
+ * Dispatch between 32-bit and 64-bit index based versions of the same algorithm
+ * implementation. This version assumes that callables for both branches consist
+ * of the same tokens, and is intended to be used with Thrust-style dispatch
+ * interfaces, that always deduce the size type from the arguments.
+ *
+ * This version of the macro supports providing two count variables, which is
+ * necessary for set algorithms.
+ */
+#define THRUST_DOUBLE_INDEX_TYPE_DISPATCH(status, call, count1, count2, arguments) \
+ if (count1 + count2 <= std::numeric_limits<thrust::detail::int32_t>::max()) { \
+ thrust::detail::int32_t THRUST_PP_CAT2(count1, _fixed) = count1; \
+ thrust::detail::int32_t THRUST_PP_CAT2(count2, _fixed) = count2; \
+ status = call arguments; \
+ } \
+ else { \
+ thrust::detail::int64_t THRUST_PP_CAT2(count1, _fixed) = count1; \
+ thrust::detail::int64_t THRUST_PP_CAT2(count2, _fixed) = count2; \
+ status = call arguments; \
+ }
+/**
+ * Dispatch between 32-bit and 64-bit index based versions of the same algorithm
+ * implementation. This version allows using different token sequences for callables
+ * in both branches, and is intended to be used with CUB-style dispatch interfaces,
+ * where the "simple" interface always forces the size to be `int` (making it harder
+ * for us to use), but the complex interface that we end up using doesn't actually
+ * provide a way to fully deduce the type from just the call, making the size type
+ * appear in the token sequence of the callable.
+ *
+ * See reduce_n_impl to see an example of how this is meant to be used.
+ */
+#define THRUST_INDEX_TYPE_DISPATCH2(status, call_32, call_64, count, arguments) \
+ if (count <= std::numeric_limits<thrust::detail::int32_t>::max()) { \
+ thrust::detail::int32_t THRUST_PP_CAT2(count, _fixed) = count; \
+ status = call_32 arguments; \
+ } \
+ else { \
+ thrust::detail::int64_t THRUST_PP_CAT2(count, _fixed) = count; \
+ status = call_64 arguments; \
+ }

MODIFY
thrust/system/hip/detail/future.inl
thrust/system/hip/detail/future.inl
@@ -918,7 +918,7 @@ public:
value_type tmp(async_signal_->extract());
async_signal_.reset();
- return std::move(tmp);
+ return tmp;
}
// For testing only.
@@ -1122,7 +1122,7 @@ void create_dependencies_impl(
{
// We only need to wait on the current dependency if we didn't steal our
// stream from it.
- if (!as.acquired_from || *as.acquired_from == I0)
+ if (!as.acquired_from || *as.acquired_from != I0)
{
create_dependency(as.stream, std::get<I0>(deps));
}

ADD
thrust/system/hip/detail/make_unsigned_special.h
+/*
+ * Copyright 2019 NVIDIA Corporation
+ * Modifications Copyright (c) 2020, Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+THRUST_BEGIN_NS
+namespace hip_rocprim {
+
+namespace detail {
+
+ template<typename Size>
+ struct make_unsigned_special;
+
+ template<>
+ struct make_unsigned_special<int> { typedef unsigned int type; };
+
+ // this is special, because HIP's atomicAdd doesn't have an overload
+ // for unsigned long, for some godforsaken reason
+ template<>
+ struct make_unsigned_special<long> { typedef unsigned long long type; };
+
+ template<>
+ struct make_unsigned_special<long long> { typedef unsigned long long type; };
+
+}
+}
+THRUST_END_NS

MODIFY
thrust/system/hip/detail/malloc_and_free.h
thrust/system/hip/detail/malloc_and_free.h
@@ -65,8 +65,8 @@ malloc(execution_policy<DerivedPolicy>&, std::size_t n)
if(status != hipSuccess)
{
- // hip_rocprim::throw_on_error(status, "device malloc failed");
- thrust::system::detail::bad_alloc(thrust::hip_category().message(status).c_str());
+ hipGetLastError(); // Clear global hip error state.
+ throw thrust::system::detail::bad_alloc(thrust::hip_category().message(status).c_str());
}
#else
result = thrust::raw_pointer_cast(thrust::malloc(thrust::seq, n));

Update system hip