HIP backend: various fixes
MODIFY
thrust/system/detail/sequential/malloc_and_free.h
thrust/system/detail/sequential/malloc_and_free.h
@@ -35,9 +35,13 @@ template<typename DerivedPolicy>
inline __host__ __device__
void *malloc(execution_policy<DerivedPolicy> &, std::size_t n)
{
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
+#if defined(__HCC__) && defined(__HIP_DEVICE_COMPILE__)
+ (void) n;
+ return 0;
+#elif !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
return std::malloc(n);
#else
+ (void) n;
return 0;
#endif
} // end mallc()
@@ -47,7 +51,9 @@ template<typename DerivedPolicy, typename Pointer>
inline __host__ __device__
void free(sequential::execution_policy<DerivedPolicy> &, Pointer ptr)
{
-#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
+#if defined(__HCC__) && defined(__HIP_DEVICE_COMPILE__)
+ (void) ptr;
+#elif !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
std::free(thrust::raw_pointer_cast(ptr));
#endif
} // end mallc()

MODIFY
thrust/system/hip/detail/error.inl
thrust/system/hip/detail/error.inl
@@ -68,7 +68,7 @@ class hip_error_category
{
using namespace hip_rocprim::errc;
- if(ev < ::hipErrorApiFailureBase)
+ if(ev < (::hipErrorMissingConfiguration - 1) /* ::hipErrorApiFailureBase */ )
{
return make_error_condition(static_cast<errc_t>(ev));
}

MODIFY
thrust/system/hip/detail/execution_policy.h
thrust/system/hip/detail/execution_policy.h
******************************************************************************/
#pragma once
+#include <hip/hip_runtime.h>
#include <thrust/detail/execution_policy.h>
#include <thrust/iterator/detail/any_system_tag.h>
#include <thrust/system/hip/config.h>

MODIFY
thrust/system/hip/detail/guarded_driver_types.h
thrust/system/hip/detail/guarded_driver_types.h
# define THRUST_DEVICE_NEEDS_RESTORATION
# endif
#else // GNUC pre 4.5.0
-# if !defined(__DRIVER_TYPES_H__)
+# if (defined(__HCC__) && !defined(HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H))
# ifdef __host__
# undef __host__
# endif
# endif // __DRIVER_TYPES_H__
#endif // __GNUC__
-
-// #include <driver_types.h>
-
+#include <hip/hcc_detail/host_defines.h>
#if !defined(__GNUC__) || ((10000 * __GNUC__ + 100 * __GNUC_MINOR__ + __GNUC_PATCHLEVEL__) >= 40500)
# ifdef THRUST_HOST_NEEDS_RESTORATION

MODIFY
thrust/system/hip/detail/guarded_hip_runtime_api.h
thrust/system/hip/detail/guarded_hip_runtime_api.h
// such as __host__ and __device__, which may already be defined by thrust
// and to undefine them before entering hip/hip_runtime_api.h (which will redefine them)
-#if defined(__HCC__) && !defined(HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H)
+#if (defined(__HCC__) && !defined(HIP_INCLUDE_HIP_HCC_DETAIL_HOST_DEFINES_H))
#ifdef __host__
#undef __host__

MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -62,7 +62,7 @@ trivial_copy_from_device(Type * dst,
size_t count,
hipStream_t stream)
{
- hipError status = hipSuccess;
+ hipError_t status = hipSuccess;
if (count == 0) return status;
status = ::hipMemcpyAsync(dst,
@@ -81,7 +81,7 @@ trivial_copy_to_device(Type * dst,
size_t count,
hipStream_t stream)
{
- hipError status = hipSuccess;
+ hipError_t status = hipSuccess;
if (count == 0) return status;
status = ::hipMemcpyAsync(dst,

MODIFY
thrust/system/hip/error.h
thrust/system/hip/error.h
@@ -177,9 +177,10 @@ using system::hip_rocprim::errc::errc_t;
} // namespace system
namespace hip
+{
// XXX replace with using system::hip_errc upon c++0x
namespace errc = system::hip::errc;
-} // end hip_rocprim
+} // end hip
using system::hip_category;

HIP backend: various fixes