Commit 7218a2b7 authored by Qianfeng Zhang's avatar Qianfeng Zhang
Browse files

Occasional tiny simplification and update in the kernel files

parent 7a7497f9
...@@ -28,13 +28,6 @@ ...@@ -28,13 +28,6 @@
// this enumerate should be synchronized with include/miopen/reduce_common.hpp // this enumerate should be synchronized with include/miopen/reduce_common.hpp
namespace ck { namespace ck {
enum class ReductionMethod_t
{
DirectThreadWise = 1,
DirectWarpWise = 2,
BlockWise = 3,
MultiBlock = 4
}; // end of namespace ck
enum class ReduceTensorOp_t enum class ReduceTensorOp_t
{ {
...@@ -71,31 +64,19 @@ enum class IndicesType_t ...@@ -71,31 +64,19 @@ enum class IndicesType_t
struct float_equal_one struct float_equal_one
{ {
template <class T>
__device__ static inline bool apply(T x)
{
return x <= type_convert<T>{}(1.0f) and x >= type_convert<T>{}(1.0f);
}
template <class T> template <class T>
__device__ inline bool operator()(T x) __device__ inline bool operator()(T x)
{ {
return (float_equal_one::apply(x)); return x <= static_cast<T>(1.0f) and x >= static_cast<T>(1.0f);
}; };
}; };
struct float_equal_zero struct float_equal_zero
{ {
template <class T>
__device__ static inline bool apply(T x)
{
return x <= type_convert<T>{}(0.0f) and x >= type_convert<T>{}(0.0f);
}
template <class T> template <class T>
__device__ inline bool operator()(T x) __device__ inline bool operator()(T x)
{ {
return (float_equal_zero::apply(x)); return x <= static_cast<T>(0.0f) and x >= static_cast<T>(0.0f);
}; };
}; };
......
...@@ -253,7 +253,7 @@ using refType_src2dDesc_padded_34 = ...@@ -253,7 +253,7 @@ using refType_src2dDesc_padded_34 =
using refType_dst1dDesc_padded = using refType_dst1dDesc_padded =
typename get_ref_desc_types<srcDims, dstDims, toReduceDims>::refType_dst1dDesc_padded; typename get_ref_desc_types<srcDims, dstDims, toReduceDims>::refType_dst1dDesc_padded;
template <ReductionMethod_t impl, bool need_padding> template <bool need_padding>
static __device__ auto get_reduction_src2d_descriptor(const void* p_src2dDesc) static __device__ auto get_reduction_src2d_descriptor(const void* p_src2dDesc)
{ {
if constexpr(need_padding) if constexpr(need_padding)
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment