Squashed 'fin/' changes from a30a51bc6..722feea66

722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
REVERT: a30a51bc6 remove unused header
REVERT: 7d2fd834c reduce scope of variable
REVERT: f6e9abe79 clang format
REVERT: 834e9a397 remove comment
REVERT: c8d6eb1a0 workspace rename
REVERT: aa7d2ea24 Merge remote-tracking branch 'origin/develop' into cderb/miopen_perf
REVERT: aaf13fb12 add to print for debug
REVERT: 34e11fa70 Merge remote-tracking branch 'origin/develop' into cderb/miopen_perf
REVERT: cb6c19d13 add search+update directives to execution context, add json examples for perf eval
REVERT: 85029077b connecting new fin functions for perf eval
REVERT: 4d1e031fd add outputs and definitions
REVERT: 952538cb8 adding perf eval function, in progress
REVERT: 617dccd9c rename
REVERT: 5c35ae886 fixes for collecting kernel blobs
REVERT: 5cfea7c43 syntax fixes
REVERT: 2f2a4ed9f add test file
REVERT: 7175019f5 first rendition of perf_compile

git-subtree-dir: fin
git-subtree-split: 722feea660e2e3d7f8e1edcc520a938be4885a44

CMakeLists.txt

@@ -39,8 +39,6 @@ else()
 endif()
 
 project ( fin C CXX )
-option( BUILD_SHARED_LIBS "Build as a shared library" ON )
-enable_testing()
 
 find_package(ROCM REQUIRED PATHS /opt/rocm)
 
@@ -53,12 +51,6 @@ include(CheckCXXCompilerFlag)
 
 option( BUILD_DEV "Build for development only" OFF)
 
-# Strip symbols for release
-if(NOT WIN32 AND NOT APPLE)
-    set(CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} -s")
-    set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s")
-endif()
-
 rocm_setup_version(VERSION 1.0.0)
 
 list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
@@ -100,7 +92,6 @@ set_property( CACHE FIN_BACKEND PROPERTY STRINGS
 
 # HIP is always required
 find_package(hip REQUIRED PATHS /opt/rocm)
-find_package(rocblas REQUIRED PATHS /opt/rocm)
 
 # OpenCL 1.2
 if( FIN_BACKEND STREQUAL "OpenCL")
@@ -327,8 +318,6 @@ else()
     message("MIOPEN_ENABLE_FIN is SET")
 endif()
 
-set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
-set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/lib)
-set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/bin)
 
 add_subdirectory(src)
+add_subdirectory(tests)

Dockerfile

@@ -1,6 +1,6 @@
 FROM ubuntu:18.04
 
-ARG PREFIX=/usr/local
+ARG PREFIX=/opt/rocm
 
 # Support multiarch
 RUN dpkg --add-architecture i386
@@ -74,21 +74,42 @@ RUN pip install https://github.com/pfultz2/cget/archive/57b3289000fcdb3b7e424c60
 # Install rclone
 RUN pip install https://github.com/pfultz2/rclone/archive/master.tar.gz
 
-# Install hcc from ROCm 3.0
-RUN rclone -b roc-3.0.x -c 286651a04d9c3a8e3052dd84b1822985498cd27d https://github.com/RadeonOpenCompute/hcc.git /hcc
-RUN LDFLAGS=-fuse-ld=gold cget -p $PREFIX install hcc,/hcc  && rm -rf /hcc
-
 # Make sure /opt/rcom is in the paths
 ENV PATH="/opt/rocm:${PATH}"
+# Install MIOpen
+ARG MIOPEN_DIR=/root/dMIOpen
+#Clone MIOpen
+RUN git clone https://github.com/ROCmSoftwarePlatform/MIOpen.git $MIOPEN_DIR
+WORKDIR $MIOPEN_DIR
+
+ARG MIOPEN_DEPS=$MIOPEN_DIR/cget
+# Install dependencies
+RUN wget --no-check-certificate -qO - https://apt.kitware.com/keys/kitware-archive-latest.asc 2>/dev/null | apt-key add -
+RUN sh -c "echo deb https://apt.kitware.com/ubuntu/ bionic main | tee -a /etc/apt/sources.list"
+RUN apt-get update && apt-get install -y --allow-unauthenticated   cmake-data=3.15.1-0kitware1 cmake=3.15.1-0kitware1
+RUN cget install pfultz2/rocm-recipes
+RUN cget install -f min-requirements.txt
+RUN CXXFLAGS='-isystem $PREFIX/include' cget install -f ./mlir-requirements.txt
 
-# Build using hcc
-RUN cget -p $PREFIX init --cxx $PREFIX/bin/hcc --std=c++14
+ARG TUNA_USER=miopenpdb
+ARG BACKEND=HIP
+# Build MIOpen
+WORKDIR $MIOPEN_DIR/build
+ARG MIOPEN_BRANCH=349928c032a42275a64a1f39c10c83daf20c1d75
+ARG MIOPEN_CACHE_DIR=/tmp/${TUNA_USER}/cache
+ARG MIOPEN_USER_DB_PATH=/tmp/$TUNA_USER/config/miopen
+ARG MIOPEN_USE_MLIR=On
+RUN git pull && git checkout $MIOPEN_BRANCH
+RUN echo "MIOPEN: Selected $BACKEND backend."
+RUN if [ $BACKEND = "OpenCL" ]; then \
+           cmake -DMIOPEN_INSTALL_CXX_HEADERS=On -DMIOPEN_INSTALL_CXX_HEADERS=On -DMIOPEN_USE_COMGR=Off -DMIOPEN_USE_MLIR=${MIOPEN_USE_MLIR} -DMIOPEN_CACHE_DIR=${MIOPEN_CACHE_DIR} -DMIOPEN_USER_DB_PATH={MIOPEN_USER_DB_PATH} -DMIOPEN_BACKEND=OpenCL -DMIOPEN_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_PREFIX_PATH="$MIOPEN_DEPS" $MIOPEN_DIR ; \
+    else \
+           CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_INSTALL_CXX_HEADERS=On -DMIOPEN_USE_COMGR=Off -DMIOPEN_USE_MLIR=${MIOPEN_USE_MLIR} -DMIOPEN_INSTALL_CXX_HEADERS=On -DMIOPEN_CACHE_DIR=${MIOPEN_CACHE_DIR} -DMIOPEN_USER_DB_PATH=${MIOPEN_USER_DB_PATH} -DMIOPEN_BACKEND=$BACKEND -DCMAKE_PREFIX_PATH=$MIOPEN_DEPS $MIOPEN_DIR ; \
+    fi
+
+RUN make -j $(nproc)
+RUN make install
 
 # Install dependencies
 ADD requirements.txt /requirements.txt
 RUN CXXFLAGS='-isystem $PREFIX/include' cget -p $PREFIX install -f /requirements.txt
-
-# Install doc requirements
-#ADD doc/requirements.txt /doc-requirements.txt
-#RUN pip install -r /doc-requirements.txt
-

Jenkinsfile

@@ -22,9 +22,8 @@ def rocmnode(name) {
 
 def cmake_build(compiler, flags, prefixpath="/opt/rocm"){
     def workspace_dir = pwd()
-    def vcache = "/var/jenkins/.cache/miopen/vcache"
     def archive = (flags == '-DCMAKE_BUILD_TYPE=release')
-    def config_targets = "all" // "check doc MIOpenDriver"
+    def config_targets = "all" 
     def test_flags = "--disable-verification-cache"
     def debug_flags = "-g -fno-omit-frame-pointer -fsanitize=undefined -fno-sanitize-recover=undefined"
     def compilerpath = ""
@@ -46,8 +45,8 @@ def cmake_build(compiler, flags, prefixpath="/opt/rocm"){
         rm -rf build
         mkdir build
         cd build
-        CXX=${compilerpath} CXXFLAGS='-Werror' cmake ${configargs} -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_FLAGS='${test_flags}' -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' ${flags} ..
-        MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS=1 CTEST_PARALLEL_LEVEL=4 MIOPEN_VERIFY_CACHE_PATH=${vcache} MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 dumb-init make -j\$(nproc) ${config_targets}
+        CXX=${compilerpath} cmake ${configargs}  ${flags} ..
+        dumb-init make -j\$(nproc) ${config_targets}
     """
     echo cmd
     sh cmd
@@ -110,26 +109,12 @@ pipeline {
         parallelsAlwaysFailFast()
     }
     environment{
-        image = "miopen"
+        image = "fin"
     }
     stages{
         // Run all static analysis tests
         stage("Static checks"){
             parallel{
-                stage('Clang Tidy') {
-                    agent{  label rocmnode("rocmtest") }
-                    environment{
-                        cmd = "rm -rf build; \
-                                mkdir build; \
-                                cd build; \
-                                CXX='clang++-3.9' cmake -DBUILD_DEV=On ..; \
-                                make -j\$(nproc) -k analyze;"
-                    }
-                    steps{
-                        buildJob('clang++-3.9', '-DCMAKE_BUILD_TYPE=release', image, "", cmd)
-                    }
-                }
-
                 stage('Clang Format') {
                     agent{ label rocmnode("rocmtest") }
                     environment{
@@ -155,11 +140,37 @@ pipeline {
                         cmd = "rm -rf build; \
                                 mkdir build; \
                                 cd build; \
-                                CXX=/usr/local/bin/hcc cmake -DBUILD_DEV=On ..; \
+                                CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On ..; \
                                 make -j\$(nproc) -k analyze;"
                     }
                     steps{
-                        buildJob('hcc', '-DCMAKE_BUILD_TYPE=release', image, "", cmd)
+                        buildJob('clang++', '-DCMAKE_BUILD_TYPE=release', image, "", cmd)
+                    }
+                }
+                stage('Build Fin') {
+                    agent{ label rocmnode("rocmtest") }
+                    environment{
+                        cmd = "rm -rf build; \
+                                mkdir build; \
+                                cd build; \
+                                CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_PREFIX_PATH=/root/dMIOpen/cget ..; \
+                                make -j\$(nproc) all;"
+                    }
+                    steps{
+                        buildJob('clang++', '-DCMAKE_BUILD_TYPE=release', image, "", cmd)
+                    }
+                }
+                stage('Fin Tests') {
+                    agent{ label rocmnode("rocmtest") }
+                    environment{
+                        cmd = "rm -rf build; \
+                                mkdir build; \
+                                cd build; \
+                                CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_PREFIX_PATH=/root/dMIOpen/cget ..; \
+                                make -j\$(nproc) all && cd tests && ctest;"
+                    }
+                    steps{
+                        buildJob('clang++', '-DCMAKE_BUILD_TYPE=release', image, "", cmd)
                     }
                 }
             }

cmake/EnableCompilerWarnings.cmake

@@ -24,7 +24,7 @@
 #
 ################################################################################
 # - Enable warning all for gcc/clang or use /W4 for visual studio
-
+function(enable_target_warnings TARGET_NAME)
 ## Strict warning level
 if (MSVC)
     # Use the highest warning level for visual studio.
@@ -105,6 +105,7 @@ else()
                 -Wno-deprecated-declarations
             )
         endif()
-        add_definitions(${CMAKE_COMPILER_WARNINGS})
+        target_compile_options(${TARGET_NAME} PRIVATE ${CMAKE_COMPILER_WARNINGS})
     endforeach()
 endif ()
+endfunction()

requirements.txt

@@ -1,5 +1,3 @@
-RadeonOpenCompute/rocm-cmake@cdd0f632b3a65bd4411593bb827eb664e25c80bc --build
-# ROCmSoftwarePlatform/MIOpen@9b2d37f9f1e4b5492ef8256cf8168363ca5fd2da 
+# RadeonOpenCompute/rocm-cmake@cdd0f632b3a65bd4411593bb827eb664e25c80bc --build
 nlohmann/json@v3.10.4
-#ROCmSoftwarePlatform/rocBLAS@9790a8658341bc665c11c311129ad0dfc533d5c4
 

src/CMakeLists.txt

@@ -34,7 +34,8 @@ configure_file("${PROJECT_SOURCE_DIR}/src/include/config.h.in" "${PROJECT_BINARY
 include_directories(include "${PROJECT_BINARY_DIR}/src/include")
 add_executable(fin main.cpp fin.cpp base64.cpp)
 target_compile_definitions( fin PRIVATE -D__HIP_PLATFORM_HCC__=1 )
-target_link_libraries(fin MIOpen ${Boost_LIBRARIES})
+enable_target_warnings(fin)
+target_link_libraries(fin MIOpen ${Boost_LIBRARIES} hip::host)
 target_link_libraries(fin ${CMAKE_THREAD_LIBS_INIT})
 if(rocblas_FOUND)
     target_link_libraries( fin $<BUILD_INTERFACE:roc::rocblas> )

src/fin.cpp

@@ -22,7 +22,7 @@ miopenDataType_t GetDataType<bfloat16>()
 {
     return miopenBFloat16;
 }
-[[gnu::noreturn]] void Fin::Usage()
+[[gnu::noreturn]] void BaseFin::Usage()
 {
     std::cout << "Usage: ./MIOpenFin *base_arg* *other_args*\n";
     std::cout << "Supported Base Arguments: conv[fp16][bfp16]\n";
@@ -37,7 +37,7 @@ void PadBufferSize(size_t& sz, int datatype_sz)
     }
 }
 
-std::string Fin::ParseBaseArg(const int argc, const char* argv[])
+std::string BaseFin::ParseBaseArg(const int argc, const char* argv[])
 {
     if(argc < 2)
     {
@@ -57,4 +57,49 @@ std::string Fin::ParseBaseArg(const int argc, const char* argv[])
     else
         return arg;
 }
+void BaseFin::InitNoGpuHandle(miopen::Handle& handle,
+                              const std::string& arch,
+                              const unsigned long num_cu)
+{
+#if MIOPEN_MODE_NOGPU
+    handle.impl->device_name        = arch;
+    handle.impl->num_cu             = num_cu;
+    handle.impl->max_mem_alloc_size = 32UL * 1024 * 1024 * 1024; // 32 GB
+    handle.impl->global_mem_size    = 32UL * 1024 * 1024 * 1024;
+    handle.impl->target_properties.Init(&handle);
+#else
+    std::ignore = handle;
+#endif
+}
+
+void BaseFin::VerifyDevProps(const std::string& in_arch, const unsigned long in_num_cu)
+{
+    std::cerr << "Verifying device properties" << std::endl;
+    std::string arch    = in_arch;
+    arch                = arch.substr(0, arch.find(':'));
+    const size_t num_cu = in_num_cu;
+    std::ignore         = num_cu;
+    if(arch == "gfx900")
+    {
+        assert(num_cu == 56 || num_cu == 64);
+    }
+    else if(arch == "gfx906")
+    {
+        assert(num_cu == 60 || num_cu == 64);
+    }
+    else if(arch == "gfx908")
+    {
+        assert(num_cu == 120);
+    }
+    else if(arch == "gfx1030")
+    {
+        assert(num_cu == 72 || num_cu == 36);
+    }
+    else if(arch == "gfx90a")
+    {
+        assert(num_cu == 110 || num_cu == 104);
+    }
+    else
+        throw std::runtime_error("Invalid Arch Name");
+}
 } // namespace fin