rjs/rk gtc

docs/sphinx/requirements.in

@@ -1 +1 @@
-rocm-docs-core[api_reference]==1.11.0
+rocm-docs-core[api_reference]==1.12.0

docs/sphinx/requirements.txt

@@ -132,7 +132,7 @@ requests==2.32.2
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core[api-reference]==1.11.0
+rocm-docs-core[api-reference]==1.12.0
     # via -r requirements.in
 six==1.16.0
     # via python-dateutil

driver/reducecalculation_driver.hpp

@@ -33,6 +33,7 @@
 #include "random.hpp"
 #include <algorithm>
 #include <cfloat>
+#include <cmath>
 #include <cstdlib>
 #include <memory>
 #include <miopen/miopen.h>
@@ -77,7 +78,7 @@ int32_t mloReduceCalculationForwardRunHost(miopenTensorDescriptor_t inputDesc,
         for(size_t i = 0; i < reduce_size; ++i)
         {
             Tcheck val = static_cast<Tcheck>(input[input_idx]);
-            if(nanPropagation && isnan(val))
+            if(nanPropagation && std::isnan(val))
             {
                 val = 0.0f;
             }

src/CMakeLists.txt

@@ -89,6 +89,7 @@ set( MIOpen_Source
     adam_api.cpp
     addlayernorm_api.cpp
     api/find2_0_commons.cpp
+    base64.cpp
     batch_norm.cpp
     batch_norm_api.cpp
     batchnorm/problem_description.cpp

src/base64.cpp

@@ -0,0 +1,155 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2024 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.
+ *
+ *******************************************************************************/
+
+#include <algorithm>
+#include <miopen/errors.hpp>
+#include <miopen/utility/base64.hpp>
+
+namespace miopen {
+
+namespace {
+
+const std::string_view base64Chars = "ABCDEFGHIJKLMNOPQRSTUVWXYZ"
+                                     "abcdefghijklmnopqrstuvwxyz"
+                                     "0123456789+/";
+
+} // namespace
+
+std::string base64Encode(const uint8_t* data, std::size_t length)
+{
+    // Calculate the exact size of the resulting Base64 string
+    size_t outputSize = ((length + 2) / 3) * 4;
+    std::string encodedString;
+    encodedString.reserve(outputSize); // Preallocate memory for the result
+
+    size_t i = 0;
+    unsigned char charArray3[3];
+    unsigned char charArray4[4];
+
+    while(length-- != 0U)
+    {
+        charArray3[i++] = *(data++);
+        if(i == 3)
+        {
+            charArray4[0] = (charArray3[0] & 0xfc) >> 2;
+            charArray4[1] = ((charArray3[0] & 0x03) << 4) + ((charArray3[1] & 0xf0) >> 4);
+            charArray4[2] = ((charArray3[1] & 0x0f) << 2) + ((charArray3[2] & 0xc0) >> 6);
+            charArray4[3] = charArray3[2] & 0x3f;
+
+            std::transform(std::begin(charArray4),
+                           std::end(charArray4),
+                           std::back_inserter(encodedString),
+                           [](unsigned char c) { return base64Chars[c]; });
+            i = 0;
+        }
+    }
+
+    if(i != 0U)
+    {
+        for(size_t j = i; j < 3; j++)
+        {
+            charArray3[j] = '\0';
+        }
+
+        charArray4[0] = (charArray3[0] & 0xfc) >> 2;
+        charArray4[1] = ((charArray3[0] & 0x03) << 4) + ((charArray3[1] & 0xf0) >> 4);
+        charArray4[2] = ((charArray3[1] & 0x0f) << 2) + ((charArray3[2] & 0xc0) >> 6);
+        charArray4[3] = charArray3[2] & 0x3f;
+
+        for(size_t j = 0; j < i + 1; j++)
+        {
+            encodedString += base64Chars[charArray4[j]];
+        }
+
+        while(i++ < 3)
+        {
+            encodedString += '=';
+        }
+    }
+
+    return encodedString;
+}
+
+std::vector<uint8_t> base64Decode(const std::string_view& encodedString)
+{
+    size_t length = encodedString.size();
+    MIOPEN_THROW_IF(length % 4 != 0, "Invalid Base64 input");
+
+    // Reserve space to avoid reallocations
+    size_t outputLength = (length / 4) * 3;
+    if(encodedString[length - 1] == '=')
+        outputLength--;
+    if(encodedString[length - 2] == '=')
+        outputLength--;
+
+    std::vector<uint8_t> decodedData;
+    decodedData.reserve(outputLength);
+
+    uint8_t charArray3[3];
+    uint8_t charArray4[4];
+    size_t i = 0;
+
+    for(char c : encodedString)
+    {
+        if(c == '=')
+            break;
+
+        auto it = std::find(base64Chars.begin(), base64Chars.end(), c);
+        MIOPEN_THROW_IF(it == base64Chars.end(), "Invalid character in Base64 string");
+
+        charArray4[i++] = static_cast<uint8_t>(std::distance(base64Chars.begin(), it));
+        if(i == 4)
+        {
+            charArray3[0] = (charArray4[0] << 2) + ((charArray4[1] & 0x30) >> 4);
+            charArray3[1] = ((charArray4[1] & 0xf) << 4) + ((charArray4[2] & 0x3c) >> 2);
+            charArray3[2] = ((charArray4[2] & 0x3) << 6) + charArray4[3];
+
+            decodedData.insert(decodedData.end(), charArray3, charArray3 + 3);
+            i = 0;
+        }
+    }
+
+    if(i != 0U)
+    {
+        for(size_t j = i; j < 4; j++)
+        {
+            charArray4[j] = 0;
+        }
+
+        charArray3[0] = (charArray4[0] << 2) + ((charArray4[1] & 0x30) >> 4);
+        charArray3[1] = ((charArray4[1] & 0xf) << 4) + ((charArray4[2] & 0x3c) >> 2);
+        charArray3[2] = ((charArray4[2] & 0x3) << 6) + charArray4[3];
+
+        for(size_t j = 0; j < i - 1; j++)
+        {
+            decodedData.push_back(charArray3[j]);
+        }
+    }
+
+    return decodedData;
+}
+
+} // namespace miopen

src/graphapi/conv_bias_res_add_activ_forward_executor.cpp

@@ -25,10 +25,11 @@
  *******************************************************************************/
 
 #include <miopen/errors.hpp>
-#include <miopen/graphapi/conv_bias_res_add_activ_forward_executor.hpp>
 #include <miopen/fusion.hpp>
+#include <miopen/graphapi/conv_bias_res_add_activ_forward_executor.hpp>
 #include <miopen/handle.hpp>
 #include <miopen/visit_float.hpp>
+#include <nlohmann/json.hpp>
 
 namespace miopen {
 
@@ -60,41 +61,73 @@ ConvolutionDescriptor Convert(const Convolution& conv, int groupCount)
 }
 } // namespace
 
+ConvBiasResAddActivForwardExecutor::ConvBiasResAddActivForwardExecutor(const nlohmann::json& json)
+    : GraphPatternExecutor(),
+      mXTensor(json.at(JsonFields::XTensor)),
+      mWTensor(json.at(JsonFields::WTensor)),
+      mConvolution(json.at(JsonFields::Convolution)),
+      mGroupCount(json.at(JsonFields::GroupCount)),
+      mZTensor(json.at(JsonFields::ZTensor)),
+      mBiasTensor(json.at(JsonFields::BiasTensor)),
+      mYTensor(json.at(JsonFields::YTensor)),
+      mAlpha1(json.at(JsonFields::Alpha1)),
+      mAlpha2(json.at(JsonFields::Alpha2)),
+      mActivationAlpha(json.at(JsonFields::ActivationAlpha))
+{
+}
+
 void ConvBiasResAddActivForwardExecutor::execute(miopenHandle_t handle, const VariantPack& vpk)
 {
-    auto convDesc = Convert(*mConvolution, mGroupCount);
+    auto convDesc = Convert(mConvolution, mGroupCount);
 
     ActivationDescriptor activDesc{miopenActivationRELU, mActivationAlpha, 1.0, 1.0};
 
-    auto* xData    = vpk.getDataPointer(mXTensor->getId());
-    auto* wData    = vpk.getDataPointer(mWTensor->getId());
-    auto* zData    = vpk.getDataPointer(mZTensor->getId());
-    auto* biasData = vpk.getDataPointer(mBiasTensor->getId());
-    auto* yData    = vpk.getDataPointer(mYTensor->getId());
+    auto* xData    = vpk.getDataPointer(mXTensor.getId());
+    auto* wData    = vpk.getDataPointer(mWTensor.getId());
+    auto* zData    = vpk.getDataPointer(mZTensor.getId());
+    auto* biasData = vpk.getDataPointer(mBiasTensor.getId());
+    auto* yData    = vpk.getDataPointer(mYTensor.getId());
 
     auto status =
         ConvBiasActivFusion(miopen::deref(handle),
                             &mAlpha1,
-                            *mXTensor,
+                            mXTensor,
                             xData,
-                            *mWTensor,
+                            mWTensor,
                             wData,
                             convDesc,
                             miopenConvFwdAlgorithm_t::miopenConvolutionFwdAlgoImplicitGEMM,
                             nullptr,
                             0,
                             &mAlpha2,
-                            *mZTensor,
+                            mZTensor,
                             zData,
-                            *mBiasTensor,
+                            mBiasTensor,
                             biasData,
                             activDesc,
-                            *mYTensor,
+                            mYTensor,
                             yData);
 
     MIOPEN_THROW_IF(status != miopenStatusSuccess, "execute failed");
 }
 
+nlohmann::json ConvBiasResAddActivForwardExecutor::getJson()
+{
+    return {
+        {GraphPatternExecutor::JsonFields::Name, name},
+        {JsonFields::XTensor, mXTensor},
+        {JsonFields::WTensor, mWTensor},
+        {JsonFields::Convolution, mConvolution},
+        {JsonFields::GroupCount, mGroupCount},
+        {JsonFields::ZTensor, mZTensor},
+        {JsonFields::BiasTensor, mBiasTensor},
+        {JsonFields::YTensor, mYTensor},
+        {JsonFields::Alpha1, mAlpha1},
+        {JsonFields::Alpha2, mAlpha2},
+        {JsonFields::ActivationAlpha, mActivationAlpha},
+    };
+}
+
 } // namespace graphapi
 
 } // namespace miopen

src/graphapi/convolution.cpp

@@ -26,6 +26,10 @@
 #include <miopen/algorithm.hpp>
 #include <miopen/graphapi/convolution.hpp>
 #include <miopen/errors.hpp>
+#include <nlohmann/json.hpp>
+
+#include <map>
+#include <string>
 
 namespace miopen {
 
@@ -1020,6 +1024,74 @@ OperationConvolution& BackendOperationConvolutionBackwardFilterDescriptor::getOp
     return mOperation;
 }
 
+void to_json(nlohmann::json& json, const Convolution& conv)
+{
+    static const std::map<miopenDataType_t, std::string> CompType2Str{
+        {miopenHalf, "miopenHalf"},
+        {miopenFloat, "miopenFloat"},
+        {miopenInt32, "miopenInt32"},
+        {miopenInt8, "miopenInt8"},
+        {miopenBFloat16, "miopenBFloat16"},
+        {miopenDouble, "miopenDouble"},
+        {miopenFloat8, "miopenFloat8"},
+        {miopenBFloat8, "miopenBFloat8"},
+        {miopenInt64, "miopenInt64"},
+    };
+
+    static const std::map<miopenConvolutionMode_t, std::string> Mode2Str{
+        {miopenConvolution, "miopenConvolution"},
+        {miopenTranspose, "miopenTranspose"},
+        {miopenGroupConv, "miopenGroupConv"},
+        {miopenDepthwise, "miopenDepthwise"},
+    };
+
+    std::string sMode     = Mode2Str.at(conv.mMode);
+    std::string sCompType = CompType2Str.at(conv.mCompType);
+
+    json = nlohmann::json{
+        {Convolution::JsonFields::SpatialDims, conv.mSpatialDims},
+        {Convolution::JsonFields::Dilations, conv.mDilations},
+        {Convolution::JsonFields::FilterStrides, conv.mFilterStrides},
+        {Convolution::JsonFields::PrePaddings, conv.mPrePaddings},
+        {Convolution::JsonFields::PostPaddings, conv.mPostPaddings},
+        {Convolution::JsonFields::CompType, sCompType},
+        {Convolution::JsonFields::Mode, sMode},
+    };
+}
+
+void from_json(const nlohmann::json& json, Convolution& conv)
+{
+    json.at(Convolution::JsonFields::SpatialDims).get_to(conv.mSpatialDims);
+    json.at(Convolution::JsonFields::Dilations).get_to(conv.mDilations);
+    json.at(Convolution::JsonFields::FilterStrides).get_to(conv.mFilterStrides);
+    json.at(Convolution::JsonFields::PrePaddings).get_to(conv.mPrePaddings);
+    json.at(Convolution::JsonFields::PostPaddings).get_to(conv.mPostPaddings);
+    auto sCompType = json.at(Convolution::JsonFields::CompType).get<std::string>();
+    auto sMode     = json.at(Convolution::JsonFields::Mode).get<std::string>();
+
+    static const std::map<std::string, miopenDataType_t> Str2CompType{
+        {"miopenHalf", miopenHalf},
+        {"miopenFloat", miopenFloat},
+        {"miopenInt32", miopenInt32},
+        {"miopenInt8", miopenInt8},
+        {"miopenBFloat16", miopenBFloat16},
+        {"miopenDouble", miopenDouble},
+        {"miopenFloat8", miopenFloat8},
+        {"miopenBFloat8", miopenBFloat8},
+        {"miopenInt64", miopenInt64},
+    };
+
+    static const std::map<std::string, miopenConvolutionMode_t> Str2Mode{
+        {"miopenConvolution", miopenConvolution},
+        {"miopenTranspose", miopenTranspose},
+        {"miopenGroupConv", miopenGroupConv},
+        {"miopenDepthwise", miopenDepthwise},
+    };
+
+    conv.mCompType = Str2CompType.at(sCompType);
+    conv.mMode     = Str2Mode.at(sMode);
+}
+
 } // namespace graphapi
 
 } // namespace miopen