HIP backend: hipify memory.h/inl
MODIFY
thrust/system/hip/detail/malloc_and_free.h
thrust/system/hip/detail/malloc_and_free.h
#pragma once
-#include <thrust/system/cuda/detail/guarded_cuda_runtime_api.h>
+#include <thrust/system/hip/detail/guarded_hip_runtime_api.h>
#include <thrust/detail/config.h>
#include <thrust/detail/seq.h>
#include <thrust/memory.h>
-#include <thrust/system/cuda/config.h>
-#ifdef THRUST_CACHING_DEVICE_MALLOC
-#include <thrust/system/cuda/detail/cub/util_allocator.cuh>
-#endif
-#include <thrust/system/cuda/detail/util.h>
+#include <thrust/system/hip/config.h>
+// STREAMHPC No caching allocator in rocPRIM
+// #ifdef THRUST_CACHING_DEVICE_MALLOC
+// #include <thrust/system/cuda/detail/cub/util_allocator.cuh>
+// #endif
+#include <thrust/system/hip/detail/util.h>
#include <thrust/system/detail/bad_alloc.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
-#ifdef THRUST_CACHING_DEVICE_MALLOC
-#define __CUB_CACHING_MALLOC
-#ifndef __CUDA_ARCH__
-inline cub::CachingDeviceAllocator &get_allocator()
-{
- static cub::CachingDeviceAllocator g_allocator(true);
- return g_allocator;
-}
-#endif
-#endif
+// STREAMHPC No caching allocator in rocPRIM
+// #ifdef THRUST_CACHING_DEVICE_MALLOC
+// #define __CUB_CACHING_MALLOC
+// #ifndef __CUDA_ARCH__
+// inline cub::CachingDeviceAllocator &get_allocator()
+// {
+// static cub::CachingDeviceAllocator g_allocator(true);
+// return g_allocator;
+// }
+// #endif
+// #endif
// note that malloc returns a raw pointer to avoid
-// depending on the heavyweight thrust/system/cuda/memory.h header
+// depending on the heavyweight thrust/system/hip/memory.h header
template<typename DerivedPolicy>
__host__ __device__
void *malloc(execution_policy<DerivedPolicy> &, std::size_t n)
{
void *result = 0;
-#ifndef __CUDA_ARCH__
-#ifdef __CUB_CACHING_MALLOC
- cub::CachingDeviceAllocator &alloc = get_allocator();
- cudaError_t status = alloc.DeviceAllocate(&result, n);
-#else
- cudaError_t status = cudaMalloc(&result, n);
-#endif
+#ifndef __HIP_DEVICE_COMPILE__
+// STREAMHPC No caching allocator in rocPRIM
+// #ifdef __CUB_CACHING_MALLOC
+// cub::CachingDeviceAllocator &alloc = get_allocator();
+// cudsError_t status = alloc.DeviceAllocate(&result, n);
+// #else
+ hipError_t status = hipMalloc(&result, n);
+// #endif
- if(status != cudaSuccess)
+ if(status != hipSuccess)
{
- // cuda_cub::throw_on_error(status, "device malloc failed");
- thrust::system::detail::bad_alloc(thrust::cuda_category().message(status).c_str());
+ // hip_rocprim::throw_on_error(status, "device malloc failed");
+ thrust::system::detail::bad_alloc(thrust::hip_category().message(status).c_str());
}
#else
result = thrust::raw_pointer_cast(thrust::malloc(thrust::seq, n));
@@ -77,18 +80,19 @@ template<typename DerivedPolicy, typename Pointer>
__host__ __device__
void free(execution_policy<DerivedPolicy> &, Pointer ptr)
{
-#ifndef __CUDA_ARCH__
-#ifdef __CUB_CACHING_MALLOC
- cub::CachingDeviceAllocator &alloc = get_allocator();
- cudaError_t status = alloc.DeviceFree(thrust::raw_pointer_cast(ptr));
-#else
- cudaError_t status = cudaFree(thrust::raw_pointer_cast(ptr));
-#endif
- cuda_cub::throw_on_error(status, "device free failed");
+#ifndef __HIP_DEVICE_COMPILE__
+// STREAMHPC No caching allocator in rocPRIM
+// #ifdef __CUB_CACHING_MALLOC
+// cub::CachingDeviceAllocator &alloc = get_allocator();
+// hipError_t status = alloc.DeviceFree(thrust::raw_pointer_cast(ptr));
+// #else
+ hipError_t status = hipFree(thrust::raw_pointer_cast(ptr));
+// #endif
+ hip_rocprim::throw_on_error(status, "device free failed");
#else
thrust::free(thrust::seq, ptr);
#endif
} // end free()
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/memory.inl
thrust/system/hip/detail/memory.inl
*/
#include <thrust/detail/config.h>
-#include <thrust/system/cuda/memory.h>
-#include <thrust/system/cuda/detail/malloc_and_free.h>
+#include <thrust/system/hip/memory.h>
+#include <thrust/system/hip/detail/malloc_and_free.h>
#include <limits>
namespace thrust
{
-// XXX WAR an issue with MSVC 2005 (cl v14.00) incorrectly implementing
-// pointer_raw_pointer for pointer by specializing it here
-// note that we specialize it here, before the use of raw_pointer_cast
-// below, which causes pointer_raw_pointer's instantiation
-#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (_MSC_VER <= 1400)
-namespace detail
-{
+// STREAMHPC Not needed for HIP
+// // XXX WAR an issue with MSVC 2005 (cl v14.00) incorrectly implementing
+// // pointer_raw_pointer for pointer by specializing it here
+// // note that we specialize it here, before the use of raw_pointer_cast
+// // below, which causes pointer_raw_pointer's instantiation
+// #if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (_MSC_VER <= 1400)
+// namespace detail
+// {
-template<typename T>
- struct pointer_raw_pointer< thrust::cuda_cub::pointer<T> >
-{
- typedef typename thrust::cuda_cub::pointer<T>::raw_pointer type;
-}; // end pointer_raw_pointer
+// template<typename T>
+// struct pointer_raw_pointer< thrust::hip_rocprim::pointer<T> >
+// {
+// typedef typename thrust::hip_rocprim::pointer<T>::raw_pointer type;
+// }; // end pointer_raw_pointer
-} // end detail
-#endif
+// } // end detail
+// #endif
-namespace cuda_cub {
+namespace hip_rocprim {
template <typename T>
template <typename OtherT>
@@ -63,25 +64,25 @@ void swap(reference<T> a, reference<T> b)
__host__ __device__
pointer<void> malloc(std::size_t n)
{
- tag cuda_tag;
- return pointer<void>(thrust::cuda_cub::malloc(cuda_tag, n));
+ tag hip_tag;
+ return pointer<void>(thrust::hip_rocprim::malloc(hip_tag, n));
} // end malloc()
template<typename T>
__host__ __device__
pointer<T> malloc(std::size_t n)
{
- pointer<void> raw_ptr = thrust::cuda_cub::malloc(sizeof(T) * n);
+ pointer<void> raw_ptr = thrust::hip_rocprim::malloc(sizeof(T) * n);
return pointer<T>(reinterpret_cast<T*>(raw_ptr.get()));
} // end malloc()
__host__ __device__
void free(pointer<void> ptr)
{
- tag cuda_tag;
- return thrust::cuda_cub::free(cuda_tag, ptr.get());
+ tag hip_tag;
+ return thrust::hip_rocprim::free(hip_tag, ptr.get());
} // end free()
-} // end cuda_
+} // end hip_rocprim
} // end thrust

MODIFY
thrust/system/hip/memory.h
thrust/system/hip/memory.h
* Copyright 2008-2013 NVIDIA Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in ccudaliance with 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
* limitations under the License.
*/
-/*! \file thrust/system/cuda/memory.h
- * \brief Managing memory associated with Thrust's CUDA system.
+/*! \file thrust/system/hip/memory.h
+ * \brief Managing memory associated with Thrust's hip system.
*/
#pragma once
#include <thrust/detail/config.h>
-#include <thrust/system/cuda/execution_policy.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>
#include <ostream>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
template <typename>
class pointer;
-} // end cuda_
+} // end hip_rocprim
END_NS_THRUST
@@ -44,10 +44,10 @@ END_NS_THRUST
BEGIN_NS_THRUST
template <typename Element>
-struct iterator_traits<thrust::cuda_cub::pointer<Element> >
+struct iterator_traits<thrust::hip_rocprim::pointer<Element> >
{
private:
- typedef thrust::cuda_cub::pointer<Element> ptr;
+ typedef thrust::hip_rocprim::pointer<Element> ptr;
public:
typedef typename ptr::iterator_category iterator_category;
@@ -57,36 +57,37 @@ public:
typedef typename ptr::reference reference;
}; // end iterator_traits
-namespace cuda_cub {
+namespace hip_rocprim {
// forward declaration of reference for pointer
template <typename Element>
class reference;
-// XXX nvcc + msvc have trouble instantiating reference below
-// this is a workaround
-template <typename Element>
-struct reference_msvc_workaround
-{
- typedef thrust::cuda_cub::reference<Element> type;
-}; // end reference_msvc_workaround
+// 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::cuda_cub::tag,
- thrust::cuda_cub::reference<T>,
- thrust::cuda_cub::pointer<T> >
+ thrust::hip_rocprim::tag,
+ thrust::hip_rocprim::reference<T>,
+ thrust::hip_rocprim::pointer<T> >
{
private:
typedef thrust::pointer<
T,
- thrust::cuda_cub::tag,
+ thrust::hip_rocprim::tag,
typename reference_msvc_workaround<T>::type,
- thrust::cuda_cub::pointer<T> >
+ thrust::hip_rocprim::pointer<T> >
super_t;
public:
@@ -124,15 +125,15 @@ template <typename T>
class reference
: public thrust::reference<
T,
- thrust::cuda_cub::pointer<T>,
- thrust::cuda_cub::reference<T> >
+ thrust::hip_rocprim::pointer<T>,
+ thrust::hip_rocprim::reference<T> >
{
private:
typedef thrust::reference<
T,
- thrust::cuda_cub::pointer<T>,
- thrust::cuda_cub::reference<T> >
+ thrust::hip_rocprim::pointer<T>,
+ thrust::hip_rocprim::reference<T> >
super_t;
public:
@@ -197,7 +198,7 @@ struct allocator
__host__ __device__ inline allocator() {}
__host__ __device__ inline allocator(const allocator &)
- : thrust::detail::malloc_allocator<T, tag, thrust::cuda_cub::pointer<T> >()
+ : thrust::detail::malloc_allocator<T, tag, thrust::hip_rocprim::pointer<T> >()
{}
template <typename U>
@@ -208,28 +209,28 @@ struct allocator
__host__ __device__ inline ~allocator() {}
}; // struct allocator
-} // namespace cuda_cub
+} // namespace hip_rocprim
namespace system {
-namespace cuda {
-using thrust::cuda_cub::pointer;
-using thrust::cuda_cub::reference;
-using thrust::cuda_cub::swap;
-using thrust::cuda_cub::malloc;
-using thrust::cuda_cub::free;
-using thrust::cuda_cub::allocator;
-} // namespace cuda
+namespace hip {
+using thrust::hip_rocprim::pointer;
+using thrust::hip_rocprim::reference;
+using thrust::hip_rocprim::swap;
+using thrust::hip_rocprim::malloc;
+using thrust::hip_rocprim::free;
+using thrust::hip_rocprim::allocator;
+} // namespace hip
} /// namespace system
-namespace cuda {
-using thrust::cuda_cub::pointer;
-using thrust::cuda_cub::reference;
-using thrust::cuda_cub::malloc;
-using thrust::cuda_cub::free;
-using thrust::cuda_cub::allocator;
-} // end cuda
+namespace hip {
+using thrust::hip_rocprim::pointer;
+using thrust::hip_rocprim::reference;
+using thrust::hip_rocprim::malloc;
+using thrust::hip_rocprim::free;
+using thrust::hip_rocprim::allocator;
+} // end hip
END_NS_THRUST
-#include <thrust/system/cuda/detail/memory.inl>
+#include <thrust/system/hip/detail/memory.inl>

HIP backend: hipify memory.h/inl