Codegen integration

CMakeLists.txt

@@ -155,7 +155,7 @@ if(NOT WIN32 AND NOT MIOPEN_WORKAROUND_USE_BOOST_FILESYSTEM)
     check_cxx_linker_flag(-lstdc++fs HAS_LIB_STD_FILESYSTEM)
 endif()
 
-list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX} ${CMAKE_INSTALL_PREFIX}/llvm ${CMAKE_INSTALL_PREFIX}/hip /opt/rocm /opt/rocm/llvm /opt/rocm/hip)
+list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX} ${CMAKE_INSTALL_PREFIX}/llvm ${CMAKE_INSTALL_PREFIX}/hip /opt/rocm /opt/rocm/llvm /opt/rocm/hip /usr/local/lib)
 
 option(ENABLE_HIP_WORKAROUNDS Off)
 set(MIOPEN_INSTALL_CXX_HEADERS Off CACHE BOOL "Install MIOpen's C++ header interface")
@@ -328,6 +328,8 @@ add_compile_definitions($<$<COMPILE_LANGUAGE:CXX>:HIP_COMPILER_FLAGS=${HIP_COMPI
 # HIP
 if( MIOPEN_BACKEND STREQUAL "HIP" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_BACKEND STREQUAL "HIPNOGPU")
     if(MIOPEN_USE_COMPOSABLEKERNEL)
+	      include(/usr/local/lib/cmake/composable_kernel_host/ck_host_targets.cmake)
+	      find_package(ck_host)
         find_package(composable_kernel 1.0.0 COMPONENTS device_other_operations device_gemm_operations device_conv_operations device_reduction_operations device_mha_operations)
     endif()
     if( MIOPEN_BACKEND STREQUAL "HIPNOGPU")

Dockerfile

@@ -114,7 +114,8 @@ DEBIAN_FRONTEND=noninteractive apt-get purge -y --allow-unauthenticated \
     miopen-hip
 
 # TODO: it should be able to automatically get commit hash from requirements.txt
-ARG CK_COMMIT=fb948120d2d674607e70d0e7587dcb249c0e74c7
+#TODO: swap my branch CK commit out with the proper commit once my branch in CK is merged
+ARG CK_COMMIT=3b9a77df7e526d2fa30534e64da3fa779c3e7d9b
 RUN wget -O ck.tar.gz https://www.github.com/ROCm/composable_kernel/archive/${CK_COMMIT}.tar.gz && \
     tar zxvf ck.tar.gz &&\
     cd composable_kernel-${CK_COMMIT} && \
@@ -125,8 +126,13 @@ RUN wget -O ck.tar.gz https://www.github.com/ROCm/composable_kernel/archive/${CK
     -D CMAKE_BUILD_TYPE=Release \
     -D GPU_ARCHS="gfx908;gfx90a;gfx942;gfx1100" \
     -D CMAKE_CXX_FLAGS=" -O3 " .. && \
-    make -j $(nproc) install 
-
+    make -j $(nproc) install && \ 
+    rm -rf CMakeCache.txt &&\ 
+    CXX=/opt/rocm/llvm/bin/clang++ cmake ../codegen \
+    -D CMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}" \
+    -D CMAKE_BUILD_TYPE=Release \
+    -D CMAKE_CXX_FLAGS=" -O3 " && \
+    make -j $(nproc) install
 # Composable Kernel installed separated from rbuild to take in values from GPU_ARCHS 
 # this can minimize build time
 RUN sed -i '/composable_kernel/d' /requirements.txt
@@ -150,4 +156,4 @@ RUN pip3 install --upgrade cmake==3.27.5
 # groupadd can add one group a time
 RUN groupadd -f render
 RUN groupadd -f video
-RUN usermod -a -G render,video root
+RUN usermod -a -G render,video root

requirements.txt

@@ -7,5 +7,5 @@ nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off
 ROCm/FunctionalPlus@v0.2.18-p0
 ROCm/eigen@3.4.0
 ROCm/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50
-ROCm/composable_kernel@fb948120d2d674607e70d0e7587dcb249c0e74c7 -DCMAKE_BUILD_TYPE=Release
+ROCm/composable_kernel@3b9a77df7e526d2fa30534e64da3fa779c3e7d9b -DCMAKE_BUILD_TYPE=Release
 google/googletest@v1.14.0

src/CMakeLists.txt

@@ -336,6 +336,7 @@ set( MIOpen_Source
     solver/softmarginloss/forward_softmarginloss.cpp
     solver/softmax/attn_softmax.cpp
     solver/softmax/softmax.cpp
+    solver/codegen_conv_hip_implicit_gemm_2d_grouped_fwd_xdlops.cpp
     subbuffers.cpp
     t5layernorm_api.cpp
     target_properties.cpp
@@ -844,10 +845,11 @@ endfunction()
 
 target_include_directories(MIOpen PUBLIC
     $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/src/include>
+    /usr/local/lib/composable_kernel_host/include
 )
 
 if(MIOPEN_USE_COMPOSABLEKERNEL)
-set(MIOPEN_CK_LINK_FLAGS composable_kernel::device_other_operations composable_kernel::device_gemm_operations composable_kernel::device_conv_operations composable_kernel::device_reduction_operations composable_kernel::device_mha_operations hip::host)
+    set(MIOPEN_CK_LINK_FLAGS composable_kernel::device_other_operations composable_kernel::device_gemm_operations composable_kernel::device_conv_operations composable_kernel::device_reduction_operations composable_kernel::device_mha_operations composable_kernel::ck_host hip::host)
 endif()
 
 if(WIN32)

src/include/miopen/conv/solvers.hpp

@@ -4494,6 +4494,26 @@ struct ConvHipImplicitGemmGroupFwdXdlops final
     bool CheckCKApplicability(const miopen::conv::ProblemDescription&) const;
 };
 
+struct ConvHipImplicitGemmGroupFwdXdlopsCodegen final : ConvSolver
+{
+    // TODO: update this fcn
+    const std::string& SolverDbId() const override
+    {
+        return GetSolverDbId<ConvHipImplicitGemmGroupFwdXdlopsCodegen>();
+    }
+
+    MIOPEN_INTERNALS_EXPORT size_t GetWorkspaceSize(
+        const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
+
+    bool MayNeedWorkspace() const override { return true; }
+
+    MIOPEN_INTERNALS_EXPORT bool
+    IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
+
+    MIOPEN_INTERNALS_EXPORT ConvSolution
+    GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
+};
+
 struct PerformanceConfigHipImplicitGemm3DGroupFwdXdlops
     : PerfConfigBaseCK<PerformanceConfigHipImplicitGemm3DGroupFwdXdlops>
 {

src/mlo_dir_conv.cpp

@@ -117,6 +117,7 @@ static auto GetImplicitGemmSolvers()
         miopen::solver::conv::ConvHipImplicitGemmGroupBwdXdlops,
         miopen::solver::conv::ConvHipImplicitGemm3DGroupFwdXdlops,
         miopen::solver::conv::ConvHipImplicitGemm3DGroupBwdXdlops,
+        miopen::solver::conv::ConvHipImplicitGemmGroupFwdXdlopsCodegen,
 #endif // MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
         miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC>{};
 }

src/solver.cpp

@@ -715,6 +715,11 @@ inline SolverRegistrar::SolverRegistrar(IdRegistryData& registry)
              multimarginloss::MultiMarginLossForward{}.SolverDbId());
 
     Register(registry, ++id, Primitive::Mha, mha::MhaCKFlashAttentionV2Forward{}.SolverDbId());
+
+    RegisterWithSolver(registry,
+                       ++id,
+                       conv::ConvHipImplicitGemmGroupFwdXdlopsCodegen{},
+                       miopenConvolutionAlgoImplicitGEMM);
     // IMPORTANT: New solvers should be added to the end of the function, and don't leave a white
     // space between this comment and the newly registered solver(s)!
 }