merge changes from branch softmax/logsoftmax optimization

src/targets/gpu/device/logsoftmax.cpp

@@ -53,7 +53,6 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
                 __syncthreads();
 
                 auto item_num = (remaining_item_num > block_size) ? block_size : remaining_item_num;
-                // reduce_max(lds_data, block_size, thr_idx, item_num, max_block_size);
                 block_reduce<type, max_op<type>>(
                     lds_data, max_op<type>{}, block_size, thr_idx, item_num, max_block_size);
 
@@ -77,7 +76,6 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
                 __syncthreads();
 
                 auto item_num = (remaining_item_num > block_size) ? block_size : remaining_item_num;
-                // reduce_sum(lds_data, block_size, thr_idx, item_num, max_block_size);
                 block_reduce<type, sum_op<type>>(
                     lds_data, sum_op<type>{}, block_size, thr_idx, item_num, max_block_size);
 

src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp

@@ -13,19 +13,19 @@ namespace device {
 template <class T>
 struct max_op
 {
-    T operator()(T x, T y) { return (x > y) ? x : y; }
+    T operator()(T x, T y) const { return (x > y) ? x : y; }
 };
 
 template <class T>
 struct min_op
 {
-    T operator()(T x, T y) { return (x < y) ? x : y; }
+    T operator()(T x, T y) const { return (x < y) ? x : y; }
 };
 
 template <class T>
 struct sum_op
 {
-    T operator()(T x, T y) { return x + y; }
+    T operator()(T x, T y) const { return x + y; }
 };
 
 template <class T, class Op>
@@ -42,7 +42,6 @@ inline __device__ void block_reduce(T* data_ptr,
         auto size   = item_num / 2;
         for(std::size_t i = thr_idx; i < size; i += block_size)
         {
-            // data_ptr[i] = ::max(to_hip_type(data_ptr[i]), to_hip_type(data_ptr[i + stride]));
             data_ptr[i] = op(data_ptr[i], data_ptr[i + stride]);
         }
         __syncthreads();
@@ -60,84 +59,6 @@ inline __device__ void block_reduce(T* data_ptr,
     __syncthreads();
 }
 
-template <class T>
-inline __device__ void reduce_argmax(T* data_ptr,
-                                     int64_t* index_ptr,
-                                     std::size_t block_size,
-                                     std::size_t thr_idx,
-                                     std::size_t item_num,
-                                     std::size_t max_index)
-{
-    while(true)
-    {
-        auto stride = (item_num + 1) / 2;
-        auto size   = item_num / 2;
-        for(std::size_t i = thr_idx; i < size; i += block_size)
-        {
-            if(data_ptr[i] < data_ptr[i + stride])
-            {
-                data_ptr[i]  = data_ptr[i + stride];
-                index_ptr[i] = index_ptr[i + stride];
-            }
-        }
-        __syncthreads();
-        item_num = stride;
-
-        if(item_num == 1)
-            break;
-    }
-
-    if(thr_idx == 0)
-    {
-        if(data_ptr[max_index] < data_ptr[0])
-        {
-            data_ptr[max_index]  = data_ptr[0];
-            index_ptr[max_index] = index_ptr[0];
-        }
-    }
-
-    __syncthreads();
-}
-
-template <class T>
-inline __device__ void reduce_argmin(T* data_ptr,
-                                     int64_t* index_ptr,
-                                     std::size_t block_size,
-                                     std::size_t thr_idx,
-                                     std::size_t item_num,
-                                     std::size_t min_index)
-{
-    while(true)
-    {
-        auto stride = (item_num + 1) / 2;
-        auto size   = item_num / 2;
-        for(std::size_t i = thr_idx; i < size; i += block_size)
-        {
-            if(data_ptr[i] > data_ptr[i + stride])
-            {
-                data_ptr[i]  = data_ptr[i + stride];
-                index_ptr[i] = index_ptr[i + stride];
-            }
-        }
-        __syncthreads();
-        item_num = stride;
-
-        if(item_num == 1)
-            break;
-    }
-
-    if(thr_idx == 0)
-    {
-        if(data_ptr[min_index] > data_ptr[0])
-        {
-            data_ptr[min_index]  = data_ptr[0];
-            index_ptr[min_index] = index_ptr[0];
-        }
-    }
-
-    __syncthreads();
-}
-
 } // namespace device
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS