Implement getitem backward

docs/reference/index.rst

@@ -33,3 +33,4 @@ The MIOpen API library is structured as follows:
   * :doc:`Cat <../doxygen/html/group__cat>` (experimental)
   * :doc:`SGD <../doxygen/html/group___s_g_d>` (experimental)
   * :doc:`ReduceExtreme <../doxygen/html/group__ReduceExtreme>` (experimental)
+  * :doc:`Getitem <../doxygen/html/group__getitem>` (experimental)

driver/CMakeLists.txt

@@ -43,6 +43,7 @@ add_executable(MIOpenDriver
     dm_dropout.cpp
     dm_fusion.cpp
     dm_gemm.cpp
+    dm_getitem.cpp
     dm_groupnorm.cpp
     dm_layernorm.cpp
     dm_lrn.cpp

driver/InputFlags.cpp

@@ -292,6 +292,165 @@ TensorParameters InputFlags::GetValueTensor(const std::string& long_name) const
 
     MIOPEN_THROW("Too many tensor descriptor parameters.");
 }
+
+TensorParametersUint64 InputFlags::GetValueTensorUint64(const std::string& long_name) const
+{
+    const auto& input     = MapInputs.at(FindShortName(long_name));
+    const auto components = miopen::SplitDelim(input.value.c_str(), ',');
+
+    if(components.size() < 1)
+        return {};
+
+    auto parse = [](auto line) {
+        auto ret        = std::vector<uint64_t>{};
+        const auto strs = miopen::SplitDelim(line, 'x');
+        for(auto&& str : strs)
+        {
+            auto elem = uint64_t{};
+            auto ss   = std::istringstream{str};
+            ss >> elem;
+
+            if(ss.bad() || ss.fail())
+                MIOPEN_THROW("Invalid tensor component " + str + " in " + line + ".");
+
+            ret.push_back(elem);
+        }
+        return ret;
+    };
+
+    auto lens = parse(components[0]);
+
+    if(components.size() == 1)
+        return {lens};
+
+    auto layout  = std::string{};
+    auto strides = std::vector<uint64_t>{};
+
+    if(std::isdigit(components[1][0]))
+        strides = parse(components[1]);
+    else
+        layout = components[1];
+
+    if(components.size() == 2)
+        return {lens, strides, layout};
+
+    MIOPEN_THROW("Too many tensor descriptor parameters.");
+}
+
+std::vector<int32_t> InputFlags::GetValueVectorInt(const std::string& long_name) const
+{
+    const auto& input = MapInputs.at(FindShortName(long_name));
+
+    auto ret        = std::vector<int32_t>{};
+    const auto strs = miopen::SplitDelim(input.value.c_str(), ',');
+
+    for(auto&& str : strs)
+    {
+        auto elem = int32_t{};
+        auto ss   = std::istringstream{str};
+        ss >> elem;
+
+        if(ss.bad() || ss.fail())
+            MIOPEN_THROW("Invalid tensor component " + str + " in " + input.value.c_str() + ".");
+
+        ret.push_back(elem);
+    }
+
+    return ret;
+}
+
+std::vector<uint64_t> InputFlags::GetValueVectorUint64(const std::string& long_name) const
+{
+    const auto& input = MapInputs.at(FindShortName(long_name));
+
+    auto ret        = std::vector<uint64_t>{};
+    const auto strs = miopen::SplitDelim(input.value.c_str(), ',');
+
+    for(auto&& str : strs)
+    {
+        auto elem = uint64_t{};
+        auto ss   = std::istringstream{str};
+        ss >> elem;
+
+        if(ss.bad() || ss.fail())
+            MIOPEN_THROW("Invalid tensor component " + str + " in " + input.value.c_str() + ".");
+
+        ret.push_back(elem);
+    }
+
+    return ret;
+}
+
+std::vector<std::vector<int32_t>>
+InputFlags::GetValue2dVectorInt(const std::string& long_name) const
+{
+    const auto& input     = MapInputs.at(FindShortName(long_name));
+    const auto components = miopen::SplitDelim(input.value.c_str(), ',');
+    auto output           = std::vector<std::vector<int32_t>>{};
+
+    if(components.size() < 1)
+        return {};
+
+    auto parse = [](auto line) {
+        auto ret        = std::vector<int32_t>{};
+        const auto strs = miopen::SplitDelim(line, 'x');
+        for(auto&& str : strs)
+        {
+            auto elem = int32_t{};
+            auto ss   = std::istringstream{str};
+            ss >> elem;
+
+            if(ss.bad() || ss.fail())
+                MIOPEN_THROW("Invalid tensor component " + str + " in " + line + ".");
+
+            ret.push_back(elem);
+        }
+        return ret;
+    };
+
+    for(auto&& component : components)
+    {
+        output.push_back(parse(component));
+    }
+
+    return output;
+}
+
+std::vector<std::vector<uint64_t>>
+InputFlags::GetValue2dVectorUint64(const std::string& long_name) const
+{
+    const auto& input     = MapInputs.at(FindShortName(long_name));
+    const auto components = miopen::SplitDelim(input.value.c_str(), ',');
+    auto output           = std::vector<std::vector<uint64_t>>{};
+
+    if(components.size() < 1)
+        return {};
+
+    auto parse = [](auto line) {
+        auto ret        = std::vector<uint64_t>{};
+        const auto strs = miopen::SplitDelim(line, 'x');
+        for(auto&& str : strs)
+        {
+            auto elem = uint64_t{};
+            auto ss   = std::istringstream{str};
+            ss >> elem;
+
+            if(ss.bad() || ss.fail())
+                MIOPEN_THROW("Invalid tensor component " + str + " in " + line + ".");
+
+            ret.push_back(elem);
+        }
+        return ret;
+    };
+
+    for(auto&& component : components)
+    {
+        output.push_back(parse(component));
+    }
+
+    return output;
+}
+
 void InputFlags::SetValue(const std::string& long_name, const std::string& new_value)
 {
     char short_name                = FindShortName(long_name);

driver/InputFlags.hpp

@@ -63,6 +63,25 @@ struct TensorParameters
     void CalculateStrides();
 };
 
+struct TensorParametersUint64
+{
+    std::vector<uint64_t> lengths = {};
+    std::vector<uint64_t> strides = {};
+    std::string layout            = "";
+
+    TensorParametersUint64 FillMissing(const TensorParametersUint64& other) const
+    {
+        return {
+            (lengths.empty() ? other.lengths : lengths),
+            (strides.empty() ? other.strides : strides),
+            (layout.empty() ? other.layout : layout),
+        };
+    }
+
+    uint64_t SetTensordDescriptor(miopenTensorDescriptor_t result, miopenDataType_t data_type);
+    void CalculateStrides();
+};
+
 class InputFlags
 {
     std::map<char, Input> MapInputs;
@@ -90,6 +109,11 @@ class InputFlags
     uint64_t GetValueUint64(const std::string& _long_name) const;
     double GetValueDouble(const std::string& _long_name) const;
     TensorParameters GetValueTensor(const std::string& long_name) const;
+    TensorParametersUint64 GetValueTensorUint64(const std::string& long_name) const;
+    std::vector<int32_t> GetValueVectorInt(const std::string& long_name) const;
+    std::vector<uint64_t> GetValueVectorUint64(const std::string& long_name) const;
+    std::vector<std::vector<int32_t>> GetValue2dVectorInt(const std::string& long_name) const;
+    std::vector<std::vector<uint64_t>> GetValue2dVectorUint64(const std::string& long_name) const;
     void SetValue(const std::string& long_name, const std::string& new_value);
     void StoreOptionalFlagValue(char short_name, const std::string& input_value);
 

driver/addlayernorm_driver.hpp

@@ -176,8 +176,8 @@ class AddLayerNormDriver : public Driver
     std::vector<Tgpu> weight;
     std::vector<Tgpu> bias;
     std::vector<Tgpu> out;
-    std::vector<Tref> mean;
-    std::vector<Tref> rstd;
+    std::vector<Tgpu> mean;
+    std::vector<Tgpu> rstd;
     std::vector<Tref> outhost;
     std::vector<Tref> meanhost;
     std::vector<Tref> rstdhost;
@@ -259,7 +259,7 @@ int AddLayerNormDriver<Tgpu, Tref>::AddCmdLineArgs()
     inflags.AddInputFlag("eps", 'e', "0.00001", "Alpha (Default=0.00001)", "double");
     inflags.AddInputFlag("normalized_dim", 'o', "3", "Nomalized Dim (Default=3)", "int");
     inflags.AddInputFlag(
-        "mode", 'm', "0", "elemwise affine mode (0), weight and bias mode (1) (Default=0)", "int");
+        "mode", 'm', "2", "elemwise affine mode (2), weight and bias mode (3) (Default=0)", "int");
 
     inflags.AddInputFlag("iter", 'i', "10", "Number of Iterations (Default=10)", "int");
     inflags.AddInputFlag("verify", 'V', "1", "Verify Each Layer (Default=1)", "int");
@@ -291,16 +291,16 @@ int AddLayerNormDriver<Tgpu, Tref>::AllocateBuffersAndCopy()
     weight_dev = std::unique_ptr<GPUMem>(new GPUMem(ctx, weight_sz, sizeof(Tgpu)));
     bias_dev   = std::unique_ptr<GPUMem>(new GPUMem(ctx, bias_sz, sizeof(Tgpu)));
     out_dev    = std::unique_ptr<GPUMem>(new GPUMem(ctx, out_sz, sizeof(Tgpu)));
-    mean_dev   = std::unique_ptr<GPUMem>(new GPUMem(ctx, mean_sz, sizeof(Tref)));
-    rstd_dev   = std::unique_ptr<GPUMem>(new GPUMem(ctx, rstd_sz, sizeof(Tref)));
+    mean_dev   = std::unique_ptr<GPUMem>(new GPUMem(ctx, mean_sz, sizeof(Tgpu)));
+    rstd_dev   = std::unique_ptr<GPUMem>(new GPUMem(ctx, rstd_sz, sizeof(Tgpu)));
 
     in       = std::vector<Tgpu>(in_sz, Tgpu0val);
     in2      = std::vector<Tgpu>(in2_sz, Tgpu0val);
     weight   = std::vector<Tgpu>(weight_sz, Tgpu0val);
     bias     = std::vector<Tgpu>(bias_sz, Tgpu0val);
     out      = std::vector<Tgpu>(out_sz, Tgpu0val);
-    mean     = std::vector<Tref>(mean_sz, Tref0val);
-    rstd     = std::vector<Tref>(rstd_sz, Tref0val);
+    mean     = std::vector<Tgpu>(mean_sz, Tgpu0val);
+    rstd     = std::vector<Tgpu>(rstd_sz, Tgpu0val);
     outhost  = std::vector<Tref>(out_sz, Tref0val);
     meanhost = std::vector<Tref>(mean_sz, Tref0val);
     rstdhost = std::vector<Tref>(rstd_sz, Tref0val);

driver/cat_driver.hpp

@@ -183,8 +183,8 @@ template <typename Tgpu, typename Tref>
 int CatDriver<Tgpu, Tref>::AddCmdLineArgs()
 {
     inflags.AddInputFlag("forw", 'F', "1", "Run only Forward Cat (Default=1)", "int");
-    inflags.AddTensorFlag("input1", '1', "", "input1 tensor descriptor");
-    inflags.AddTensorFlag("input2", '2', "", "input2 tensor descriptor");
+    inflags.AddTensorFlag("input1", '1', "2x32x128x128x128", "input1 tensor descriptor");
+    inflags.AddTensorFlag("input2", '2', "2x32x128x128x128", "input2 tensor descriptor");
     inflags.AddTensorFlag("input3", '3', "", "input3 tensor descriptor");
     inflags.AddTensorFlag("input4", '4', "", "input4 tensor descriptor");
     inflags.AddTensorFlag("input5", '5', "", "input5 tensor descriptor");

driver/dm_getitem.cpp

@@ -0,0 +1,40 @@
+/*******************************************************************************
+ *
+ * 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 "getitem_driver.hpp"
+#include "registry_driver_maker.hpp"
+
+static Driver* makeDriver(const std::string& base_arg)
+{
+    if(base_arg == "getitem")
+        return new GetitemDriver<float, float>();
+    if(base_arg == "getitemfp16")
+        return new GetitemDriver<float16, float>();
+    if(base_arg == "getitembfp16")
+        return new GetitemDriver<bfloat16, float>();
+    return nullptr;
+}
+
+REGISTER_DRIVER_MAKER(makeDriver);

driver/driver.hpp

@@ -169,13 +169,13 @@ inline void PadBufferSize(size_t& sz, int datatype_sz)
 [[noreturn]] inline void Usage()
 {
     printf("Usage: ./driver *base_arg* *other_args*\n");
-    printf("Supported Base Arguments: conv[fp16|int8|bfp16|fp8|bfp8], CBAInfer[fp16], "
-           "pool[fp16], lrn[fp16], "
+    printf("Supported Base Arguments: conv[fp16|int8|bfp16], pool[fp16], lrn[fp16], "
            "activ[fp16], softmax[fp16], bnorm[fp16], rnn[fp16], gemm[fp16], ctc, dropout[fp16], "
-           "tensorop[fp16], reduce[fp16|fp64], layernorm[bfp16|fp16], sum[bfp16|fp16], "
+           "tensorop, reduce[fp16|fp64], layernorm[bfp16|fp16], sum[bfp16|fp16], "
            "groupnorm[bfp16|fp16], cat[bfp16|fp16], addlayernorm[bfp16|fp16], "
            "t5layernorm[bfp16|fp16], adam[fp16], ampadam, reduceextreme[bfp16|fp16], "
-           "adamw[fp16], ampadamw, transformersadamw[fp16], transformersampadamw\n");
+           "adamw[fp16], ampadamw, transformersadamw[fp16], transformersampadamw, "
+           "getitem[bfp16|fp16]\n");
     exit(0); // NOLINT (concurrency-mt-unsafe)
 }
 
@@ -190,22 +190,22 @@ inline std::string ParseBaseArg(int argc, char* argv[])
     std::string arg = argv[1];
 
     if(arg != "conv" && arg != "convfp16" && arg != "convint8" && arg != "convbfp16" &&
-       arg != "convfp8" && arg != "convbfp8" && arg != "CBAInfer" && arg != "CBAInferfp16" &&
        arg != "pool" && arg != "poolfp16" && arg != "lrn" && arg != "lrnfp16" && arg != "activ" &&
        arg != "activfp16" && arg != "softmax" && arg != "softmaxfp16" && arg != "bnorm" &&
        arg != "bnormfp16" && arg != "rnn" && arg != "rnnfp16" && arg != "rnn_seq" &&
        arg != "rnn_seqfp16" && arg != "gemm" && arg != "gemmfp16" && arg != "ctc" &&
-       arg != "dropout" && arg != "dropoutfp16" && arg != "tensorop" && arg != "tensoropfp16" &&
-       arg != "reduce" && arg != "reducefp16" && arg != "reducefp64" && arg != "layernorm" &&
-       arg != "layernormfp16" && arg != "layernormbfp16" && arg != "sum" && arg != "sumfp16" &&
-       arg != "sumbfp16" && arg != "groupnorm" && arg != "groupnormfp16" &&
-       arg != "groupnormbfp16" && arg != "cat" && arg != "catfp16" && arg != "catbfp16" &&
-       arg != "addlayernorm" && arg != "addlayernormfp16" && arg != "addlayernormbfp16" &&
-       arg != "t5layernorm" && arg != "t5layernormfp16" && arg != "t5layernormbfp16" &&
-       arg != "adam" && arg != "adamfp16" && arg != "ampadam" && arg != "reduceextreme" &&
+       arg != "dropout" && arg != "dropoutfp16" && arg != "tensorop" && arg != "reduce" &&
+       arg != "reducefp16" && arg != "reducefp64" && arg != "layernorm" && arg != "layernormfp16" &&
+       arg != "layernormbfp16" && arg != "sum" && arg != "sumfp16" && arg != "sumbfp16" &&
+       arg != "groupnorm" && arg != "groupnormfp16" && arg != "groupnormbfp16" && arg != "cat" &&
+       arg != "catfp16" && arg != "catbfp16" && arg != "addlayernorm" &&
+       arg != "addlayernormfp16" && arg != "addlayernormbfp16" && arg != "t5layernorm" &&
+       arg != "t5layernormfp16" && arg != "t5layernormbfp16" && arg != "adam" &&
+       arg != "adamfp16" && arg != "ampadam" && arg != "reduceextreme" &&
        arg != "reduceextremefp16" && arg != "reduceextremebfp16" && arg != "adamw" &&
        arg != "adamwfp16" && arg != "ampadamw" && arg != "transformersadamw" &&
-       arg != "transformersadamwfp16" && arg != "transformersampadamw" && arg != "--version")
+       arg != "transformersadamwfp16" && arg != "transformersampadamw" && arg != "getitem" &&
+       arg != "getitemfp16" && arg != "getitembfp16" && arg != "--version")
     {
         printf("FAILED: Invalid Base Input Argument\n");
         Usage();