Remove constexpr from initialized zeroVal and tiny fix in reduction_operator.hpp

composable_kernel/include/tensor_operation/gridwise_generic_2d_reduction_blockwise.hpp

@@ -92,7 +92,7 @@ struct GridwiseReduction_xy_to_x_blockwise
         // LDS
         __shared__ compType p_in_block_buffer[BlockBufferSize];
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -243,7 +243,7 @@ struct GridwiseReduction_xy_to_x_blockwise
         __shared__ compType p_in_block_buffer[BlockBufferSize];
         __shared__ int block_indices_buffer[BlockBufferSize];
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -431,7 +431,7 @@ struct GridwiseReduction_xy_to_x_blockwise
         __shared__ compType p_in_block_buffer[BlockBufferSize];
         __shared__ int block_indices_buffer[BlockBufferSize];
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_val_buf =
             make_dynamic_buffer<AddressSpaceEnum_t::Global>(ws_values_global,

composable_kernel/include/tensor_operation/gridwise_generic_2d_reduction_direct_threadwise.hpp

@@ -82,7 +82,7 @@ struct GridwiseReduction_xy_to_x_direct_threadwise
         (void)ws_indices_global;
         (void)indices_global;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -204,7 +204,7 @@ struct GridwiseReduction_xy_to_x_direct_threadwise
     {
         (void)ws_indices_global;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -348,7 +348,7 @@ struct GridwiseReduction_xy_to_x_direct_threadwise
     {
         (void)origReduceLen;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_val_buf =
             make_dynamic_buffer<AddressSpaceEnum_t::Global>(ws_values_global,

composable_kernel/include/tensor_operation/gridwise_generic_2d_reduction_direct_warpwise.hpp

@@ -82,7 +82,7 @@ struct GridwiseReduction_xy_to_x_direct_warpwise
         (void)ws_indices_global;
         (void)indices_global;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -215,7 +215,7 @@ struct GridwiseReduction_xy_to_x_direct_warpwise
     {
         (void)ws_indices_global;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_buf = make_dynamic_buffer<AddressSpaceEnum_t::Global>(
             p_src_global, src2dDesc.GetElementSpaceSize(), type_convert<srcDataType>{}(zeroVal));
@@ -373,7 +373,7 @@ struct GridwiseReduction_xy_to_x_direct_warpwise
     {
         (void)origReduceLen;
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         const auto src_global_val_buf =
             make_dynamic_buffer<AddressSpaceEnum_t::Global>(ws_values_global,

composable_kernel/include/tensor_operation/gridwise_generic_2d_reduction_multiblock.hpp

@@ -86,7 +86,7 @@ struct GridwiseReduction_xy_to_x_multiblock
         (void)alpha; // unused
         (void)beta;  // unused
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         // LDS
         __shared__ compType p_in_block_buffer[BlockBufferSize];
@@ -216,7 +216,7 @@ struct GridwiseReduction_xy_to_x_multiblock
         (void)alpha; // unused
         (void)beta;  // unused
 
-        constexpr auto zeroVal = opReduce::GetReductionZeroVal();
+        const auto zeroVal = opReduce::GetReductionZeroVal();
 
         // LDS
         __shared__ compType p_in_block_values_buffer[BlockBufferSize];

composable_kernel/include/utility/reduction_operator.hpp

@@ -82,7 +82,7 @@ struct Max
 {
     using dataType = T;
 
-    __device__ static constexpr T GetReductionZeroVal() { return NumericLimits<T>::lowest(); };
+    __device__ static constexpr T GetReductionZeroVal() { return NumericLimits<T>::Lowest(); };
 
     __device__ inline constexpr void operator()(T& a, T b) const
     {