AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'amd-develop' into amd-master

Browse Source 
Change-Id: I05572d2b32f1df70b54e2efeb32c8a4d8055912d
(cherry picked from commit 3a56e5c09b)


[ROCm/clr commit: cc4a73f215]
This commit is contained in:
Maneesh Gupta
2017-04-13 12:39:28 +05:30
committed by Maneesh Gupta
parent 6b9e7127ad
commit 0f0f7c8e61
14 changed files with 374 additions and 129 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/grid_launch_GGL.hpp
+20 -13
View File
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
#pragma once
#if GENERIC_GRID_LAUNCH == 1
#include "concepts.hpp"
@@ -71,7 +72,7 @@ namespace hip_impl
template<FunctionalProcedure F, typename... Ts>
using is_new_grid_launch_t = typename std::conditional<
std::is_callable<F(Ts...)>{},
is_callable<F(Ts...)>{},
New_grid_launch_tag,
Old_grid_launch_tag>::type;
}
@@ -118,6 +119,7 @@ namespace hip_impl
// TODO: these are workarounds, they should be removed.
hc::accelerator_view lock_stream_hip_(hipStream_t&, void*&);
void print_prelaunch_trace_(const char*, dim3, dim3, int, hipStream_t);
void unlock_stream_hip_(
hipStream_t, void*, const char*, hc::accelerator_view*);
@@ -137,7 +139,13 @@ namespace hip_impl
void* lck_stream = nullptr;
auto acc_v = lock_stream_hip_(stream, lck_stream);
auto stream_guard = make_RAII_guard(
[](){ /* perhaps use a slimmed down ihipPrintKernelLaunch here */ },
std::bind(
print_prelaunch_trace_,
kernel_name,
num_blocks,
dim_blocks,
group_mem_bytes,
stream),
std::bind(
unlock_stream_hip_, stream, lck_stream, kernel_name, &acc_v));
@@ -841,16 +849,15 @@ namespace hip_impl
group_mem_bytes,\
stream,\
...)\
do {\
hipLaunchKernelGGL(\
kernel_name,\
num_blocks,\
dim_blocks,\
group_mem_bytes,\
stream,\
hipLaunchParm{},\
##__VA_ARGS__);\
} while(0)
do {\
hipLaunchKernelGGL(\
kernel_name,\
num_blocks,\
dim_blocks,\
group_mem_bytes,\
stream,\
hipLaunchParm{},\
##__VA_ARGS__);\
} while(0)
}
#endif //GENERIC_GRID_LAUNCH
projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp
+70 -80
View File
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
#pragma once
#include "concepts.hpp"
#include <type_traits> // For std::conditional, std::decay, std::enable_if,
// std::false_type, std result_of and std::true_type.
@@ -29,9 +30,6 @@ THE SOFTWARE.
namespace std
{ // TODO: these should be removed as soon as possible.
#if (__cplusplus < 201406L)
template<typename...>
using void_t = void;
#if (__cplusplus < 201402L)
template<bool cond, typename T = void>
using enable_if_t = typename enable_if<cond, T>::type;
@@ -43,88 +41,80 @@ namespace std
using result_of_t = typename result_of<F(Ts...)>::type;
template<typename T>
using remove_reference_t = typename remove_reference<T>::type;
template<
FunctionalProcedure F,
unsigned int n = 0u,
typename = void>
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
// Pointer to member function, call through non-pointer.
template<FunctionalProcedure F, typename C, typename... Ts>
struct is_callable_impl<
F(C, Ts...),
0u,
void_t<decltype((declval<C>().*declval<F>())(declval<Ts>()...))>
> : true_type {
};
// Pointer to member function, call through pointer.
template<FunctionalProcedure F, typename C, typename... Ts>
struct is_callable_impl<
F(C, Ts...),
1u,
void_t<decltype(((*declval<C>()).*declval<F>())(declval<Ts>()...))>
> : std::true_type {
};
// Pointer to member data, call through non-pointer, no args.
template<FunctionalProcedure F, typename C>
struct is_callable_impl<
F(C),
2u,
void_t<decltype(declval<C>().*declval<F>())>
> : true_type {
};
// Pointer to member data, call through pointer, no args.
template<FunctionalProcedure F, typename C>
struct is_callable_impl<
F(C),
3u,
void_t<decltype(*declval<C>().*declval<F>())>
> : true_type {
};
// General call, n args.
template<FunctionalProcedure F, typename... Ts>
struct is_callable_impl<
F(Ts...),
4u,
void_t<decltype(declval<F>()(declval<Ts>()...))>
> : true_type {
};
// Not callable.
template<FunctionalProcedure F>
struct is_callable_impl<F, 5u> : false_type {};
template<typename Call>
struct is_callable : is_callable_impl<Call> {};
#else
template<typename, typename = void>
struct is_callable_impl : false_type {};
template<FunctionalProcedure F, typename... Ts>
struct is_callable_impl<
F(Ts...),
void_t<result_of_t<F(Ts...)>>> : true_type {};
template<typename F>
struct is_callable : is_callable_impl<F> {};
#endif
template<typename...>
struct disjunction : false_type {};
template<typename B1>
struct disjunction<B1> : B1 {};
template<typename B1, typename... Bs>
struct disjunction<B1, Bs...>
: conditional_t<B1{} == true, B1, disjunction<Bs...>>
{};
#endif
}
namespace hip_impl // Only for documentation, macros ignore namespaces.
namespace hip_impl
{
template<typename...>
using void_t_ = void;
#if (__cplusplus < 201402L)
template<
FunctionalProcedure F,
unsigned int n = 0u,
typename = void>
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
// Pointer to member function, call through non-pointer.
template<FunctionalProcedure F, typename C, typename... Ts>
struct is_callable_impl<
F(C, Ts...),
0u,
void_t_<decltype((std::declval<C>().*std::declval<F>())(
std::declval<Ts>()...))>
> : std::true_type {};
// Pointer to member function, call through pointer.
template<FunctionalProcedure F, typename C, typename... Ts>
struct is_callable_impl<
F(C, Ts...),
1u,
void_t_<decltype(((*std::declval<C>()).*std::declval<F>())(
std::declval<Ts>()...))>
> : std::true_type {};
// Pointer to member data, call through non-pointer, no args.
template<FunctionalProcedure F, typename C>
struct is_callable_impl<
F(C),
2u,
void_t_<decltype(std::declval<C>().*std::declval<F>())>
> : std::true_type {};
// Pointer to member data, call through pointer, no args.
template<FunctionalProcedure F, typename C>
struct is_callable_impl<
F(C),
3u,
void_t_<decltype(*std::declval<C>().*std::declval<F>())>
> : std::true_type {};
// General call, n args.
template<FunctionalProcedure F, typename... Ts>
struct is_callable_impl<
F(Ts...),
4u,
void_t_<decltype(std::declval<F>()(std::declval<Ts>()...))>
> : std::true_type {};
// Not callable.
template<FunctionalProcedure F>
struct is_callable_impl<F, 5u> : std::false_type {};
template<typename Call>
struct is_callable : is_callable_impl<Call> {};
#else
template<typename, typename = void>
struct is_callable_impl : std::false_type {};
template<FunctionalProcedure F, typename... Ts>
struct is_callable_impl<
F(Ts...),
void_t_<std::result_of_t<F(Ts...)>>> : std::true_type {};
#endif
#define count_macro_args_impl_hip_(\
_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\
_16, _17, _18, _19, _20, _21, _22, _23, _24, _25, _26, _27, _28, _29,\
projects/clr/hipamd/include/hip/hcc_detail/hip_complex.h
+1
View File
@@ -24,6 +24,7 @@ THE SOFTWARE.
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
#include "hip/hcc_detail/hip_vector_types.h"
#include <math.h>
#if __cplusplus
#define COMPLEX_ADD_OP_OVERLOAD(type) \
projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
+37 -1
View File
@@ -1270,6 +1270,15 @@ __device__ __host__ static inline type operator op (type& val, int) { \
#define DECLOP_1VAR_COMP(type, op) \
__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
return lhs.x op rhs.x; \
} \
__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \
return lhs.x op rhs.x; \
} \
__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \
return lhs.x op rhs.x ; \
} \
__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \
return lhs.x op rhs.x ; \
}
#define DECLOP_1VAR_1IN_1OUT(type, op) \
@@ -1338,6 +1347,15 @@ __device__ __host__ static inline type operator op (type& val, int) { \
#define DECLOP_2VAR_COMP(type, op) \
__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y); \
} \
__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y); \
} \
__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y); \
} \
__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y); \
}
#define DECLOP_2VAR_1IN_1OUT(type, op) \
@@ -1415,7 +1433,16 @@ __device__ __host__ static inline type operator op (type& val, int) { \
#define DECLOP_3VAR_COMP(type, op) \
__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \
}
} \
__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \
} \
__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \
} \
__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \
} \
#define DECLOP_3VAR_1IN_1OUT(type, op) \
__device__ __host__ static inline type operator op(type &rhs) { \
@@ -1500,6 +1527,15 @@ __device__ __host__ static inline type operator op (type& val, int) { \
#define DECLOP_4VAR_COMP(type, op) \
__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \
} \
__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \
} \
__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \
} \
__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \
return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \
}
#define DECLOP_4VAR_1IN_1OUT(type, op) \
projects/clr/hipamd/include/hip/hcc_detail/host_defines.h
+1 -1
View File
@@ -48,7 +48,7 @@ THE SOFTWARE.
#define __global__ __attribute__((hc_grid_launch)) __attribute__((used))
#else
//#warning "GGL global define reached"
#define __global__ [[hc]] __attribute__((weak))
#define __global__ __attribute__((hc, weak))
#endif //GENERIC_GRID_LAUNCH
#define __noinline__ __attribute__((noinline))
projects/clr/hipamd/include/hip/hcc_detail/math_functions.h
+1
View File
@@ -51,6 +51,7 @@ __device__ float exp10f(float x);
__device__ float exp2f(float x);
__device__ float expf(float x);
__device__ float expm1f(float x);
__device__ int abs(int x);
__device__ float fabsf(float x);
__device__ float fdimf(float x, float y);
__device__ float fdividef(float x, float y);
projects/clr/hipamd/include/hip/hip_fp16.h
+1 -1
View File
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#ifdef HIP_INCLUDE_HIP_HIP_FP16_H
#ifndef HIP_INCLUDE_HIP_HIP_FP16_H
#define HIP_INCLUDE_HIP_HIP_FP16_H
#include <hip/hip_common.h>