diff --git a/cmake/rocm.cmake b/cmake/rocm.cmake
index e5f49b116ac76cf91f356dc56b1fb2a647b6a722..6a7cb00814293242f41caf8d41c46c8da94341cd 100644
--- a/cmake/rocm.cmake
+++ b/cmake/rocm.cmake
@@ -13,89 +13,79 @@ else()
     message(FATAL_ERROR "Could not find HIP. Ensure that HIP is either installed in /opt/rocm/hip or the variable HIP_PATH is set to point to the right location.")
 endif()
 
-string(REPLACE "." ";" HIP_VERSION_LIST ${HIP_VERSION})
-list(GET HIP_VERSION_LIST 0 HIP_VERSION_MAJOR)
-list(GET HIP_VERSION_LIST 1 HIP_VERSION_MINOR)
-if (NOT ${HIP_VERSION_MAJOR} STREQUAL "3")
-    message(FATAL_ERROR "ROCM version needed 3.x, Please update ROCM.")
-else()
-    if (${HIP_VERSION_MINOR} LESS "7")
-        message(WARNING "ROCM version 3.x which x(got ${HIP_VERSION_MINOR}) greater equal 7 is prefered.")
+if (${HIP_VERSION} VERSION_LESS 3.0)
+    message(FATAL_ERROR "ROCM version needed 3. Please update ROCM.")
+endif()
+
+macro(hipconfig_get_option variable option)
+    if(NOT DEFINED ${variable})
+        execute_process(
+            COMMAND ${HIP_HIPCONFIG_EXECUTABLE} ${option}
+            OUTPUT_VARIABLE ${variable})
     endif()
+endmacro()
+
+hipconfig_get_option(HIP_COMPILER "--compiler")
+hipconfig_get_option(HIP_CPP_CONFIG "--cpp_config")
+
+separate_arguments(HIP_CPP_CONFIG)
+
+foreach(hip_config_item ${HIP_CPP_CONFIG})
+    foreach(macro_name "__HIP_PLATFORM_HCC__" "__HIP_ROCclr__")
+        if(${hip_config_item} STREQUAL "-D${macro_name}=")
+            set(HIP_CPP_DEFINE "${HIP_CPP_DEFINE}#define ${macro_name}\n")
+            set(HIP_CPP_UNDEFINE "${HIP_CPP_UNDEFINE}\
+                    #ifdef ${macro_name}\n#undef ${macro_name}\n\
+                    #else\n#error\n\
+                    #endif\n")
+        elseif(${hip_config_item} STREQUAL "-D${macro_name}")
+            set(HIP_CPP_DEFINE "${HIP_CPP_DEFINE}#define ${macro_name} 1\n")
+            set(HIP_CPP_UNDEFINE "${HIP_CPP_UNDEFINE}\
+                    #ifdef ${macro_name}\n#undef ${macro_name}\n\
+                    #else\n#error\n\
+                    #endif\n")
+        endif()
+    endforeach()
+endforeach()
+
+message(STATUS "Using HIP compiler ${HIP_COMPILER}")
+
+if(${HIP_COMPILER} STREQUAL "hcc")
+    set(MGE_ROCM_LIBS hip_hcc)
+    message(WARNING "hcc is not well supported, please modify link.txt to link with hipcc")
+elseif (${HIP_COMPILER} STREQUAL "clang")
+    set(MGE_ROCM_LIBS amdhip64)
 endif()
 
-set(MGE_ROCM_LIBS OpenCL amdhip64 MIOpen rocblas rocrand)
+list(APPEND MGE_ROCM_LIBS amdocl64 MIOpen rocblas rocrand)
 
 set(HIP_INCLUDE_DIR ${HIP_ROOT_DIR}/../include)
 set(HIP_LIBRARY_DIR ${HIP_ROOT_DIR}/../lib)
 
-#miopen
-get_filename_component(__found_miopen_library ${HIP_ROOT_DIR}/../miopen/lib REALPATH)
-find_path(MIOPEN_LIBRARY_DIR
-    NAMES libMIOpen.so
-    HINTS ${PC_MIOPEN_INCLUDE_DIRS} ${MIOPEN_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_miopen_library}
-    PATH_SUFFIXES lib
-    DOC "Path to MIOPEN library directory." )
-
-if(MIOPEN_LIBRARY_DIR STREQUAL "MIOPEN_LIBRARY_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find MIOPEN Library")
-endif()
-
-get_filename_component(__found_miopen_include ${HIP_ROOT_DIR}/../miopen/include REALPATH)
-find_path(MIOPEN_INCLUDE_DIR
-    NAMES miopen
-    HINTS ${PC_MIOPEN_INCLUDE_DIRS} ${MIOPEN_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_miopen_include}
-    PATH_SUFFIXES include
-    DOC "Path to MIOPEN include directory." )
-
-if(MIOPEN_INCLUDE_DIR STREQUAL "MIOPEN_INCLUDE_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find MIOEPN INCLUDE")
-endif()
-
-#rocblas
-get_filename_component(__found_rocblas_library ${HIP_ROOT_DIR}/../rocblas/lib REALPATH)
-find_path(ROCBLAS_LIBRARY_DIR
-    NAMES librocblas.so
-    HINTS ${PC_ROCBLAS_INCLUDE_DIRS} ${ROCBLAS_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_rocblas_library}
-    PATH_SUFFIXES lib
-    DOC "Path to ROCBLAS library directory." )
-
-if(ROCBLAS_LIBRARY_DIR STREQUAL "ROCBLAS_LIBRARY_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find ROCBLAS Library")
-endif()
-
-get_filename_component(__found_rocblas_include ${HIP_ROOT_DIR}/../rocblas/include REALPATH)
-find_path(ROCBLAS_INCLUDE_DIR
-    NAMES rocblas.h
-    HINTS ${PC_ROCBLAS_INCLUDE_DIRS} ${ROCBLAS_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_rocblas_include}
-    PATH_SUFFIXES include
-    DOC "Path to ROCBLAS include directory." )
-
-if(ROCBLAS_INCLUDE_DIR STREQUAL "ROCBLAS_INCLUDE_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find ROCBLAS INCLUDE")
-endif()
-
-#rocrand
-get_filename_component(__found_rocrand_library ${HIP_ROOT_DIR}/../rocrand/lib REALPATH)
-find_path(ROCRAND_LIBRARY_DIR
-    NAMES librocrand.so
-    HINTS ${PC_ROCRAND_INCLUDE_DIRS} ${ROCRAND_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_rocrand_library}
-    PATH_SUFFIXES lib
-    DOC "Path to ROCRAND library directory." )
+function(find_rocm_library name dirname include library)
+    find_path(${name}_LIBRARY_DIR
+        NAMES ${library}
+        HINTS "${HIP_ROOT_DIR}/../${dirname}"
+        PATH_SUFFIXES lib lib/x86_64
+        DOC "Path to ${name} library directory")
 
-if(ROCRAND_LIBRARY_DIR STREQUAL "ROCRAND_LIBRARY_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find ROCRAND Library")
-endif()
+    if(${${name}_LIBRARY_DIR} MATCHES "NOTFOUND$")
+        message(FATAL_ERROR "Can not find ${name} library")
+    endif()
 
-get_filename_component(__found_rocrand_include ${HIP_ROOT_DIR}/../rocrand/include REALPATH)
-find_path(ROCRAND_INCLUDE_DIR
-    NAMES rocrand.h
-    HINTS ${PC_ROCRAND_INCLUDE_DIRS} ${ROCRAND_ROOT_DIR} ${ROCM_TOOLKIT_INCLUDE} ${__found_rocrand_include}
-    PATH_SUFFIXES include
-    DOC "Path to ROCRAND include directory." )
-
-if(ROCRAND_INCLUDE_DIR STREQUAL "ROCRAND_INCLUDE_DIR-NOTFOUND")
-    message(FATAL_ERROR "Can not find ROCRAND INCLUDE")
-endif()
+    find_path(${name}_INCLUDE_DIR
+        NAMES ${include}
+        HINTS "${HIP_ROOT_DIR}/../${dirname}"
+        PATH_SUFFIXES include
+        DOC "Path to ${name} include directory")
 
+    if(${name}_INCLUDE_DIR MATCHES "NOTFOUND$")
+        message(FATAL_ERROR "Can not find ${name} include")
+    endif()
+    message(DEBUG "Found lib ${${name}_LIBRARY_DIR}, include ${${name}_INCLUDE_DIR}")
+endfunction()
 
+find_rocm_library(MIOPEN miopen miopen libMIOpen.so)
+find_rocm_library(ROCBLAS rocblas rocblas.h librocblas.so)
+find_rocm_library(ROCRAND rocrand rocrand.h librocrand.so)
+find_rocm_library(AMDOCL opencl CL libamdocl64.so)
diff --git a/dnn/include/hcc_detail/hcc_defs_epilogue.h b/dnn/include/hcc_detail/hcc_defs_epilogue.h.in
similarity index 76%
rename from dnn/include/hcc_detail/hcc_defs_epilogue.h
rename to dnn/include/hcc_detail/hcc_defs_epilogue.h.in
index bb20636093c6a016dd8c0b2d9480659bcb8b4c20..5baa616e3f8c636bd7c0a222fcf69026b0daeb2a 100644
--- a/dnn/include/hcc_detail/hcc_defs_epilogue.h
+++ b/dnn/include/hcc_detail/hcc_defs_epilogue.h.in
@@ -9,10 +9,6 @@
  * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  */
 
-#ifdef __HIP_PLATFORM_HCC__
-#undef __HIP_PLATFORM_HCC__
-#else
-#error "hcc_defs_epilogue.h must be included after hcc_defs_prologue.h"
-#endif
+@HIP_CPP_UNDEFINE@
 
 // vim: syntax=cpp.doxygen
diff --git a/dnn/include/hcc_detail/hcc_defs_prologue.h b/dnn/include/hcc_detail/hcc_defs_prologue.h.in
similarity index 93%
rename from dnn/include/hcc_detail/hcc_defs_prologue.h
rename to dnn/include/hcc_detail/hcc_defs_prologue.h.in
index e93eb410ccc80739ac70cc104e753fc69038d038..e6e9c043d5196d7001fd1ef6ebab49efb40c7381 100644
--- a/dnn/include/hcc_detail/hcc_defs_prologue.h
+++ b/dnn/include/hcc_detail/hcc_defs_prologue.h.in
@@ -9,6 +9,6 @@
  * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  */
 
-#define __HIP_PLATFORM_HCC__
+@HIP_CPP_DEFINE@
 
 // vim: syntax=cpp.doxygen
diff --git a/dnn/src/CMakeLists.txt b/dnn/src/CMakeLists.txt
index 25e4cfe1d60685dcc1b4af87d9737c92557c5c68..c1990c4b1a848fc28c4ddba3bc5794f02028a276 100644
--- a/dnn/src/CMakeLists.txt
+++ b/dnn/src/CMakeLists.txt
@@ -63,7 +63,7 @@ macro (HIP_COMPILE _hip_target _hip_objs)
     add_custom_target(${_hip_target})
 
     # set return value
-    set (${_hip_objs} ${_generated_files})
+    set(${_hip_objs} ${_generated_files})
 endmacro()
 
 if (MGE_WITH_ROCM)
@@ -74,14 +74,21 @@ if (MGE_WITH_ROCM)
     # empty file to bypass this error.
     file(GLOB start.cpp.hip "" )
     list(APPEND HIP_SOURCES start.cpp.hip)
+    configure_file(
+        ${PROJECT_SOURCE_DIR}/dnn/include/hcc_detail/hcc_defs_prologue.h.in
+        ${PROJECT_BINARY_DIR}/dnn/include/hcc_detail/hcc_defs_prologue.h)
 
-    file (GLOB_RECURSE HIPSOURCES rocm/*.cpp.hip)
-    set(HIP_TARGET_NAME hip_kernel)
+    configure_file(
+        ${PROJECT_SOURCE_DIR}/dnn/include/hcc_detail/hcc_defs_epilogue.h.in
+        ${PROJECT_BINARY_DIR}/dnn/include/hcc_detail/hcc_defs_epilogue.h)
+
+    file(GLOB_RECURSE HIP_SOURCES_ rocm/*.cpp.hip)
+    set(HIP_TARGET_NAME megdnn_hip_kernel)
     set(_HIPCC_OPTIONS "-fPIC")
     set(_HCC_OPTIONS "-fPIC")
     set(_NVCC_OPTIONS "-fPIC")
 
-    list(APPEND HIP_SOURCES ${HIPSOURCES})
+    list(APPEND HIP_SOURCES ${HIP_SOURCES_})
     set_source_files_properties(${HIP_SOURCES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
     HIP_INCLUDE_DIRECTORIES(${PROJECT_SOURCE_DIR}/dnn
         ${PROJECT_SOURCE_DIR}/dnn/include
@@ -91,13 +98,14 @@ if (MGE_WITH_ROCM)
         ${HIP_INCLUDE_DIR}
         ${MIOPEN_INCLUDE_DIR}
         ${ROCBLAS_INCLUDE_DIR}
-        ${ROCRAND_INCLUDE_DIR})
+        ${ROCRAND_INCLUDE_DIR}
+        ${AMDOCL_INCLUDE_DIR})
     hip_compile(
-       ${HIP_TARGET_NAME} HIPOBJS ${HIP_SOURCES}
-       HIPCC_OPTIONS ${_HIPCC_OPTIONS}
-       HCC_OPTIONS ${_HCC_OPTIONS}
-       NVCC_OPTIONS ${_NVCC_OPTIONS})
-    list (APPEND SOURCES ${HIPOBJS})
+        ${HIP_TARGET_NAME} HIPOBJS ${HIP_SOURCES}
+        HIPCC_OPTIONS ${_HIPCC_OPTIONS}
+        HCC_OPTIONS ${_HCC_OPTIONS}
+        NVCC_OPTIONS ${_NVCC_OPTIONS})
+    list(APPEND SOURCES ${HIPOBJS})
 endif ()
 
 if(MGE_WITH_CUDA)
@@ -139,16 +147,18 @@ if(MGE_WITH_CUDA)
 endif()
 
 if(MGE_WITH_ROCM)
-        target_include_directories(megdnn PUBLIC
+    target_include_directories(megdnn PUBLIC
         ${HIP_INCLUDE_DIR}
         ${MIOPEN_INCLUDE_DIR}
         ${ROCBLAS_INCLUDE_DIR}
-        ${ROCRAND_INCLUDE_DIR})
-        target_link_directories(megdnn PUBLIC
+        ${ROCRAND_INCLUDE_DIR}
+        ${AMDOCL_INCLUDE_DIR})
+    target_link_directories(megdnn PUBLIC
         ${HIP_LIBRARY_DIR}
         ${MIOPEN_LIBRARY_DIR}
         ${ROCBLAS_LIBRARY_DIR}
-        ${ROCRAND_LIBRARY_DIR})
+        ${ROCRAND_LIBRARY_DIR}
+        ${AMDOCL_LIBRARY_DIR})
 endif()
 
 
diff --git a/dnn/src/fallback/batched_matrix_mul/opr_impl.cpp b/dnn/src/fallback/batched_matrix_mul/opr_impl.cpp
index a9a190748703015683f01375c70d39287cc88170..758a6da8ce9b6026ecd2254196096f16563f6550 100644
--- a/dnn/src/fallback/batched_matrix_mul/opr_impl.cpp
+++ b/dnn/src/fallback/batched_matrix_mul/opr_impl.cpp
@@ -11,7 +11,6 @@
  */
 #include "./opr_impl.h"
 #include "./algos.h"
-#include "hcc_detail/hcc_defs_prologue.h"
 
 #include "src/common/algo_chooser.h"
 #include "src/common/utils.cuh"
diff --git a/dnn/src/rocm/miopen_with_check.h b/dnn/src/rocm/miopen_with_check.h
index 24d8eca8309df275055cd243e9c4a058c73c2cec..a7ebed3635cac6ad393466d5b9f21e7df5d8ca6d 100644
--- a/dnn/src/rocm/miopen_with_check.h
+++ b/dnn/src/rocm/miopen_with_check.h
@@ -11,9 +11,7 @@
 
 #pragma once
 
-#ifndef __HIP_PLATFORM_HCC__
-#define __HIP_PLATFORM_HCC__
-#endif
+#include "hcc_detail/hcc_defs_prologue.h"
 
 #include <miopen/version.h>
 #pragma GCC diagnostic push