Merge branch 'develop' into asm_igemm_nhwc_fwd_bwd

# RESOLVED Conflicts:
#	src/solver.cpp

src/CMakeLists.txt

@@ -115,6 +115,8 @@ set( MIOpen_Source
     execution_context.cpp
     reducetensor.cpp
     reducetensor_api.cpp
+    activ/problem_description.cpp
+    solver/activ/fwd_0.cpp
     include/miopen/buffer_info.hpp
     include/miopen/temp_file.hpp
     include/miopen/bfloat16.hpp

src/activ/problem_description.cpp

@@ -0,0 +1,75 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 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 <miopen/activ/problem_description.hpp>
+#include <miopen/names.hpp>
+
+#include <sstream>
+
+namespace miopen {
+
+namespace activ {
+
+NetworkConfig ProblemDescription::MakeNetworkConfig() const
+{
+    // short cut for packed tensors and 2D tensors with stride != width
+    const auto& x_lens = xDesc.GetLengths();
+
+    const auto x_elem_sz = xDesc.GetElementSize();
+
+    const auto x_width2D =
+        ((x_lens.size() == 2) ? x_lens[1] : (x_lens.size() == 3) ? x_lens[2] : (x_lens.size() == 4)
+                                                                                   ? x_lens[3]
+                                                                                   : x_lens[4]);
+
+    const auto height =
+        (x_lens.size() == 2) ? x_lens[0] : (x_lens.size() == 3) ? x_lens[1] : (x_lens.size() == 4)
+                                                                                  ? x_lens[2]
+                                                                                  : x_lens[3];
+
+    const auto packed = xDesc.IsPacked() && yDesc.IsPacked();
+
+    const auto read_len = (packed) ? x_elem_sz : x_width2D;
+
+    const auto read_unit = (read_len % 4 == 0) ? 4 : (read_len % 2 == 0) ? 2 : 1;
+    const auto MAP_RD    = read_len / read_unit;
+
+    std::ostringstream ss;
+
+    ss << "activ-";
+    ss << ((packed) ? "11" : "10"); // + lite bit
+    ss << xDesc.GetType();
+    ss << activDesc.GetMode();
+    ss << read_unit;
+    ss << MAP_RD;
+    ss << height;
+
+    return NetworkConfig{ss.str()};
+}
+
+} // namespace activ
+
+} // namespace miopen

src/binary_cache.cpp

@@ -60,17 +60,15 @@ static boost::filesystem::path ComputeSysCachePath()
 static boost::filesystem::path ComputeUserCachePath()
 {
 #ifdef MIOPEN_CACHE_DIR
-    std::string cache_dir = MIOPEN_CACHE_DIR;
-    std::string version;
-
-    version = std::to_string(MIOPEN_VERSION_MAJOR) + "." + std::to_string(MIOPEN_VERSION_MINOR) +
-              "." + std::to_string(MIOPEN_VERSION_PATCH) + "." +
-              MIOPEN_STRINGIZE(MIOPEN_VERSION_TWEAK);
-    auto p = boost::filesystem::path{miopen::ExpandUser(cache_dir)} / version;
+    const std::string cache_dir = MIOPEN_CACHE_DIR;
+    const std::string version =
+        std::to_string(MIOPEN_VERSION_MAJOR) + "." + std::to_string(MIOPEN_VERSION_MINOR) + "." +
+        std::to_string(MIOPEN_VERSION_PATCH) + "." + MIOPEN_STRINGIZE(MIOPEN_VERSION_TWEAK);
 
     const char* const custom = miopen::GetStringEnv(MIOPEN_CUSTOM_CACHE_DIR{});
-    if(custom != nullptr && strlen(custom) > 0)
-        p = boost::filesystem::path{miopen::ExpandUser(custom)};
+    const auto p             = (custom != nullptr && strlen(custom) > 0)
+                       ? boost::filesystem::path{miopen::ExpandUser(custom)}
+                       : boost::filesystem::path{miopen::ExpandUser(cache_dir)} / version;
 
     if(!boost::filesystem::exists(p) && !MIOPEN_DISABLE_USERDB)
         boost::filesystem::create_directories(p);

src/expanduser.cpp

@@ -7,7 +7,7 @@ MIOPEN_DECLARE_ENV_VAR(HOME)
 
 namespace miopen {
 
-std::string ExpandUser(std::string p)
+std::string ExpandUser(const std::string& p)
 {
     const char* home_dir = GetStringEnv(HOME{});
     if(home_dir == nullptr || home_dir == std::string("/") || home_dir == std::string(""))

src/hip/handlehip.cpp

@@ -503,12 +503,11 @@ std::size_t Handle::GetGlobalMemorySize() const
 
 std::size_t Handle::GetMaxComputeUnits() const
 {
+    const std::size_t num_cu = Value(MIOPEN_DEVICE_CU{});
+    if(num_cu > 0)
+        return num_cu;
+
     int result;
-    const char* const num_cu = miopen::GetStringEnv(MIOPEN_DEVICE_CU{});
-    if(num_cu != nullptr && strlen(num_cu) > 0)
-    {
-        return boost::lexical_cast<std::size_t>(num_cu);
-    }
     auto status =
         hipDeviceGetAttribute(&result, hipDeviceAttributeMultiprocessorCount, this->impl->device);
     if(status != hipSuccess)

src/hipoc/hipoc_program.cpp

@@ -58,8 +58,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_OPENCL_ENFORCE_CODE_OBJECT_OPTION)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_OPENCL_ENFORCE_CODE_OBJECT_VERSION)
 MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEVICE_ARCH)
 
-#define MIOPEN_WORKAROUND_SWDEV_225285 1
-
 #if MIOPEN_USE_COMGR
 #define MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_27 1
 #endif
@@ -164,19 +162,12 @@ static hipModulePtr CreateModule(const boost::filesystem::path& hsaco_file)
 template <typename T> /// intended for std::string and std::vector<char>
 hipModulePtr CreateModuleInMem(const T& blob)
 {
-#if !MIOPEN_WORKAROUND_SWDEV_225285
     hipModule_t raw_m;
     auto status = hipModuleLoadData(&raw_m, reinterpret_cast<const void*>(blob.data()));
     hipModulePtr m{raw_m};
     if(status != hipSuccess)
         MIOPEN_THROW_HIP_STATUS(status, "Failed loading module");
     return m;
-#else
-    TmpDir tmp_dir("miopen");
-    auto file_path = tmp_dir.path / boost::filesystem::unique_path("miopen-%%%%-%%%%-%%%%-%%%%");
-    WriteFile(blob, file_path);
-    return CreateModule(file_path);
-#endif
 }
 
 HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name,
@@ -187,16 +178,12 @@ HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name,
 }
 
 HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name, const std::string& blob)
-    : program(program_name)
+    : program(program_name) ///, module(CreateModuleInMem(blob))
 {
-    TmpDir tmp_dir("miopen");
-    auto file_path = tmp_dir.path / boost::filesystem::unique_path("miopen-%%%%-%%%%-%%%%-%%%%");
-    WriteFile(blob, file_path);
-    const char* const arch = miopen::GetStringEnv(MIOPEN_DEVICE_ARCH{});
-    if(arch == nullptr)
-    {
-        this->module = CreateModule(file_path);
-    }
+    if(nullptr !=
+       miopen::GetStringEnv(MIOPEN_DEVICE_ARCH{})) /// \todo Finish off this spaghetti eventually.
+        return;
+    module = CreateModuleInMem(blob);
 }
 
 HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name,

src/include/miopen/activ/invoke_params.hpp

@@ -0,0 +1,52 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 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 activ {
+
+struct InvokeParams : public miopen::InvokeParams
+{
+    InvokeParams() = default;
+
+    double alpha = 0;
+    TensorDescriptor x_desc;
+    ConstData_t x = nullptr;
+    double beta   = 0;
+    TensorDescriptor y_desc;
+    Data_t y        = nullptr;
+    double gamma    = 0;
+    size_t x_offset = 0;
+    size_t y_offset = 0;
+};
+
+} // namespace activ
+
+} // namespace miopen

src/include/miopen/activ/problem_description.hpp

@@ -0,0 +1,80 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 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/activ.hpp>
+#include <miopen/tensor.hpp>
+
+#include <string>
+
+namespace miopen {
+
+struct NetworkConfig;
+
+namespace activ {
+
+enum class Direction
+{
+    Forward,
+    Backward,
+};
+
+struct ProblemDescription
+{
+    ProblemDescription(Direction direction_,
+                       const ActivationDescriptor& activ,
+                       const TensorDescriptor& xDesc_,
+                       const TensorDescriptor& yDesc_)
+        : direction(direction_), activDesc(activ), xDesc(xDesc_), yDesc(yDesc_)
+    {
+    }
+
+    Direction GetDirection() const { return direction; }
+    const ActivationDescriptor& GetActivDesc() const { return activDesc; }
+    const TensorDescriptor& GetXDesc() const { return xDesc; }
+    const TensorDescriptor& GetYDesc() const { return yDesc; }
+
+    NetworkConfig MakeNetworkConfig() const;
+
+    void Serialize(std::ostream& stream) const;
+
+    friend std::ostream& operator<<(std::ostream& os, const ProblemDescription& obj)
+    {
+        obj.Serialize(os);
+        return os;
+    }
+
+    private:
+    Direction direction;
+    ActivationDescriptor activDesc;
+    TensorDescriptor xDesc;
+    TensorDescriptor yDesc;
+};
+
+} // namespace activ
+
+} // namespace miopen

src/include/miopen/activ/solvers.hpp

@@ -0,0 +1,53 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 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>
+
+namespace miopen {
+
+namespace activ {
+struct ProblemDescription;
+} // namespace activ
+
+namespace solver {
+
+namespace activ {
+
+struct ActivFwdSolver0 : public SolverBase<ProblemDescription>
+{
+    bool IsApplicable(const ExecutionContext& context,
+                      const miopen::activ::ProblemDescription& problem) const;
+    ConvSolution GetSolution(const ExecutionContext& context,
+                             const miopen::activ::ProblemDescription& problem) const;
+};
+
+} // namespace activ
+
+} // namespace solver
+
+} // namespace miopen

src/include/miopen/execution_context.hpp

@@ -84,6 +84,7 @@ struct ExecutionContext
     inline void SetStream(Handle* stream_) { stream = stream_; }
 
     ExecutionContext() = default;
+    ExecutionContext(Handle* stream_) : stream(stream_) {}
 
     void DetectRocm();
 

src/include/miopen/expanduser.hpp

@@ -30,7 +30,7 @@
 
 namespace miopen {
 
-std::string ExpandUser(std::string p);
+std::string ExpandUser(const std::string& p);
 
 } // namespace miopen