Enable device-side memory allocation on HIP
However, at this moment device-side memory allocation in HIP
is under development and is not usable (any kernels which use
malloc/free crash with mem access fault).
MODIFY
thrust/system/detail/sequential/malloc_and_free.h
thrust/system/detail/sequential/malloc_and_free.h
@@ -37,8 +37,7 @@ inline __host__ __device__
void *malloc(execution_policy<DerivedPolicy> &, std::size_t n)
{
#if defined(__HIP_DEVICE_COMPILE__)
- (void) n;
- return 0;
+ return ::malloc(n);
#elif !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
return std::malloc(n);
#else
@@ -53,7 +52,7 @@ inline __host__ __device__
void free(sequential::execution_policy<DerivedPolicy> &, Pointer ptr)
{
#if defined(__HIP_DEVICE_COMPILE__)
- (void) ptr;
+ ::free(thrust::raw_pointer_cast(ptr));
#elif !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 200)
std::free(thrust::raw_pointer_cast(ptr));
#endif

MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -104,7 +104,7 @@ trivial_copy_device_to_device(Policy& policy, Type* dst, Type const* src, size_t
inline void __host__ __device__ terminate()
{
#ifdef __HIP_DEVICE_COMPILE__
- asm("trap;");
+ abort();
#else
std::terminate();
#endif

MODIFY
thrust/system/hip/memory.h
thrust/system/hip/memory.h
#include <hip/hip_runtime.h>
#include <ostream>
-#include <thrust/detail/allocator/malloc_allocator.h>
#include <thrust/detail/config.h>
-#include <thrust/detail/type_traits.h>
-#include <thrust/memory.h>
#include <thrust/system/hip/execution_policy.h>
+#include <thrust/memory.h>
+#include <thrust/detail/type_traits.h>
+#include <thrust/detail/allocator/malloc_allocator.h>
BEGIN_NS_THRUST
namespace hip_rocprim
@@ -67,15 +67,6 @@ namespace hip_rocprim
template <typename Element>
class reference;
-// STREAMHPC Not needed for HIP
-// // XXX nvcc + msvc have trouble instantiating reference below
-// // this is a workaround
-// template <typename Element>
-// struct reference_msvc_workaround
-// {
-// typedef thrust::hip_rocprim::reference<Element> type;
-// }; // end reference_msvc_workaround
-
template <typename T>
class pointer : public thrust::pointer<T,
thrust::hip_rocprim::tag,
@@ -86,7 +77,7 @@ class pointer : public thrust::pointer<T,
private:
typedef thrust::pointer<T,
thrust::hip_rocprim::tag,
- typename reference_msvc_workaround<T>::type,
+ thrust::hip_rocprim::reference<T>,
thrust::hip_rocprim::pointer<T>>
super_t;

Enable device-side memory allocation on HIP

However, at this moment device-side memory allocation in HIP
is under development and is not usable (any kernels which use
malloc/free crash with mem access fault).