Softmax ocl refactoring

src/CMakeLists.txt

@@ -144,7 +144,9 @@ set( MIOpen_Source
     rnn_api.cpp
     rnn/rnn_util.cpp
     rnn/Solutions/rnn_transformer.cpp
+    softmax.cpp
     softmax_api.cpp
+    softmax/problem_description.cpp
     solution.cpp
     solver.cpp
     solver/activ/bwd_0.cpp
@@ -249,6 +251,7 @@ set( MIOpen_Source
     solver/pooling/backward2d.cpp
     solver/pooling/backwardNd.cpp
     solver/reduce/forward_sum.cpp
+    solver/softmax/softmax.cpp
     subbuffers.cpp
     sum_api.cpp
     target_properties.cpp
@@ -560,7 +563,6 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
         ocl/mloNorm.cpp
         ocl/pooling_ocl.cpp
         ocl/tensorocl.cpp
-        ocl/softmaxocl.cpp
         ocl/rnnocl.cpp
         ocl/utilocl.cpp
         ocl/ctcocl.cpp

src/include/miopen/ctc.hpp

@@ -56,7 +56,7 @@ struct CTCLossDescriptor : miopenCTCLossDescriptor
                                    const int* inputLengths,
                                    miopenCTCLossAlgo_t algo) const;
 
-    void CTCLoss(const Handle& handle,
+    void CTCLoss(Handle& handle,
                  const TensorDescriptor& probsDesc,
                  ConstData_t probs,
                  const int* labels,

src/include/miopen/softmax.hpp

@@ -34,7 +34,7 @@ namespace miopen {
 struct Handle;
 struct TensorDescriptor;
 
-miopenStatus_t SoftmaxForward(const Handle& handle,
+miopenStatus_t SoftmaxForward(Handle& handle,
                               const void* alpha,
                               const void* beta,
                               const TensorDescriptor& xDesc,

src/include/miopen/softmax/invoke_params.hpp

@@ -0,0 +1,68 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#pragma once
+
+#include <miopen/invoke_params.hpp>
+#include <miopen/tensor.hpp>
+
+namespace miopen {
+namespace softmax {
+
+struct InvokeParams : public miopen::InvokeParams
+{
+    InvokeParams( const void* alpha_,
+                            const void* beta_,
+                            const TensorDescriptor& xDesc_,
+                            ConstData_t x_,
+                            const TensorDescriptor& yDesc_,
+                            Data_t y_,
+                            miopenSoftmaxAlgorithm_t algorithm_,
+                            miopenSoftmaxMode_t mode_,
+                            int x_offset_ = 0,
+                            int y_offset_ = 0)
+        : alpha(alpha_), beta(beta_), xDesc(xDesc_), x(x_), yDesc(yDesc_), y(y_), algorithm(algorithm_), mode(mode_), x_offset(x_offset_), y_offset(y_offset_)
+    {
+    }
+
+    std::size_t GetWorkspaceSize() const { return 0; }
+    Data_t GetWorkspace() const { return nullptr; }
+
+    //miopenHandle_t handle;
+    const void* alpha;
+    const void* beta;
+    const TensorDescriptor& xDesc;
+    ConstData_t x;
+    const TensorDescriptor& yDesc;
+    Data_t y;
+    miopenSoftmaxAlgorithm_t algorithm;
+    miopenSoftmaxMode_t mode;
+    int x_offset = 0;
+    int y_offset = 0;
+};
+
+} // namespace cat
+} // namespace miopen

src/include/miopen/softmax/problem_description.hpp

@@ -0,0 +1,90 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#pragma once
+
+#include <miopen/problem_description_base.hpp>
+#include <miopen/tensor.hpp>
+
+
+namespace miopen {
+
+struct NetworkConfig;
+
+namespace softmax {
+
+int nextPow2(int v);
+
+void getParams(const TensorDescriptor& in_yDesc, 
+                miopenSoftmaxMode_t in_mode, 
+                int& out_n, 
+                int& out_c, 
+                int& out_h, 
+                int& out_w, 
+                int& out_grid_size, 
+                int& out_spatial_dim,
+                int& out_vector_size,
+                int& out_num_batch,
+                bool& out_usefp16,
+                bool& out_usefp32,
+                std::vector<size_t>& out_vld,
+                std::vector<size_t>& out_vgd,
+                size_t& out_workgroups,
+                int& out_batch_size,
+                int& out_u_batch_size);
+
+struct ProblemDescription : ProblemDescriptionBase
+{
+    ProblemDescription( const void* alpha_,
+                        const void* beta_,
+                        const TensorDescriptor& xDesc_,
+                        const TensorDescriptor& yDesc_,
+                        miopenSoftmaxAlgorithm_t algorithm_,
+                        miopenSoftmaxMode_t mode_) :
+        alpha(alpha_), beta(beta_), xDesc(xDesc_), yDesc(yDesc_), algorithm(algorithm_), mode(mode_)
+    {
+    }
+
+    const void* GetAlpha() const {return alpha;}
+    const void* GetBeta() const {return beta;}
+    const TensorDescriptor& GetXDesc() const {return xDesc;}
+    const TensorDescriptor& GetYDesc() const {return yDesc;}
+    const miopenSoftmaxAlgorithm_t GetAlgorithm() const {return algorithm;}
+    const miopenSoftmaxMode_t GetMode() const {return mode;}
+
+    NetworkConfig MakeNetworkConfig() const override;    
+
+private:
+    const void* alpha;
+    const void* beta;
+    const TensorDescriptor& xDesc;
+    const TensorDescriptor& yDesc;
+    const miopenSoftmaxAlgorithm_t algorithm;
+    const miopenSoftmaxMode_t mode;
+};
+
+} // namespace softmax
+} // namespace miopen

src/include/miopen/softmax/solvers.hpp

@@ -0,0 +1,65 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2023 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#pragma once
+
+#include <miopen/solver.hpp>
+#include <miopen/softmax/problem_description.hpp>
+
+#include <utility>
+
+namespace miopen {
+
+namespace solver {
+
+namespace softmax {
+
+using SoftmaxSolver = NonTunableSolverBase<ExecutionContext, miopen::softmax::ProblemDescription>;
+
+struct SoftmaxForward final : SoftmaxSolver
+{
+    const std::string& SolverDbId() const override { return GetSolverDbId<SoftmaxForward>(); }
+
+    bool IsApplicable(const ExecutionContext& context,
+                      const miopen::softmax::ProblemDescription& problem) const override;
+
+    ConvSolution GetSolution(const ExecutionContext& context,
+                             const miopen::softmax::ProblemDescription& problem) const override;
+
+    std::size_t GetWorkspaceSize(const ExecutionContext& context,
+                                 const miopen::softmax::ProblemDescription& problem) const override
+    {
+        return 0;
+    }
+
+    bool MayNeedWorkspace() const override { return false; }
+};
+
+} // namespace cat
+
+} // namespace solver
+
+} // namespace miopen

src/ocl/ctcocl.cpp

@@ -44,7 +44,7 @@
 
 namespace miopen {
 
-void CTCLossDescriptor::CTCLoss(const Handle& handle,
+void CTCLossDescriptor::CTCLoss(Handle& handle,
                                 const TensorDescriptor& probsDesc,
                                 ConstData_t probs,
                                 const int* labels,