Merge pull request #1259 from zaliu/master

ROCm 4.1 merge staging into master

.gitattributes

@@ -0,0 +1,9 @@
+/CHANGELOG.md merge=union
+*.pdf binary
+*.doc binary
+*.docx binary
+*.ppt binary
+*.pptx binary
+*.xls binary
+*.xlsx binary
+*.xps binary

.githooks/pre-commit

@@ -29,6 +29,7 @@ fi
 
 # Change the copyright date at the top of any text files
 for file in $files; do
+    [[ -L $file ]] && continue
     echo "Processing copyright dates in $file"
     if [[ -e $file ]]; then
         /usr/bin/perl -pi -e 'INIT { exit 1 if !-f $ARGV[0] || -B $ARGV[0]; $year = (localtime)[5] + 1900 }
@@ -39,6 +40,7 @@ done
 
 # do the formatting
 for file in $files; do
+    [[ -L $file ]] && continue
     if [[ -e $file ]] && echo $file | grep -Eq '\.c$|\.h$|\.hpp$|\.cpp$|\.cl$|\.in$|\.txt$|\.yaml$|\.yml$|\.sh$|\.py$|\.pl$|\.cmake$|\.md$|\.rst$|\.groovy$|\.ini$|\.awk$|\.csv$'; then
         echo "Processing line endings in $file"
         sed -i -e 's/[[:space:]]*$//' "$file" # Remove whitespace at end of lines
@@ -50,6 +52,7 @@ done
 # if clang-format exists, run it on C/C++ files
 if command -v clang-format >/dev/null; then
     for file in $files; do
+       [[ -L $file ]] && continue
        if [[ -e $file ]] && echo $file | grep -Eq '\.c$|\.h$|\.hpp$|\.cpp$|\.cl$|\.h\.in$|\.hpp\.in$|\.cpp\.in$'; then
            echo "clang-format $file"
            clang-format -i -style=file "$file"

.gitignore

@@ -6,6 +6,9 @@ Tensile.egg-info
 build*
 dist
 .cache
+/*.yaml
+/*.txt
+/*log*
 
 0_Build
 1_BenchmarkProblems

.jenkins/common.groovy

@@ -62,7 +62,7 @@ def runCompileCommand(platform, project, jobName, boolean debug=false)
     junit "${project.paths.project_build_prefix}/python_unit_tests.xml"
 }
 
-def publishResults(project)
+def publishResults(project, boolean skipHostTest=false)
 {
     try
     {
@@ -72,7 +72,7 @@ def publishResults(project)
     {
         try
         {
-            junit "${project.paths.project_build_prefix}/build/host_test_output.xml"
+            if (!skipHostTest) junit "${project.paths.project_build_prefix}/build/host_test_output.xml"
         }
         finally
         {
@@ -81,12 +81,13 @@ def publishResults(project)
     }
 }
 
-def runTestCommand (platform, project, jobName, test_marks)
+def runTestCommand (platform, project, jobName, test_marks, boolean skipHostTest=false)
 {
     def test_dir =  "Tensile/Tests"
 
     String compiler = jobName.contains('hipclang') ? 'hipcc' : 'hcc'
     String pythonVersion = 'py36'
+    String markSkipHostTest = skipHostTest ? "#" : ""
 
     def command = """#!/usr/bin/env bash
             set -x
@@ -98,11 +99,11 @@ def runTestCommand (platform, project, jobName, test_marks)
 
             gpuArch=`/opt/rocm/bin/rocm_agent_enumerator  | tail -n 1`
 
-            pushd build
-            ./TensileTests --gtest_output=xml:host_test_output.xml --gtest_color=yes
-            HOST_ERR=\$?
+            ${markSkipHostTest}pushd build
+            ${markSkipHostTest}./TensileTests --gtest_output=xml:host_test_output.xml --gtest_color=yes
+            ${markSkipHostTest}HOST_ERR=\$?
+            ${markSkipHostTest}popd
 
-            popd
             #### temporary fix to remedy incorrect home directory
             export HOME=/home/jenkins
             ####
@@ -112,10 +113,10 @@ def runTestCommand (platform, project, jobName, test_marks)
             PY_ERR=\$?
             date
 
-            if [[ \$HOST_ERR -ne 0 ]]
-            then
-                exit \$HOST_ERR
-            fi
+            ${markSkipHostTest}if [[ \$HOST_ERR -ne 0 ]]
+            ${markSkipHostTest}then
+            ${markSkipHostTest}    exit \$HOST_ERR
+            ${markSkipHostTest}fi
 
             if [[ \$PY_ERR -ne 0 ]]
             then
@@ -134,14 +135,14 @@ def runTestCommand (platform, project, jobName, test_marks)
     {
         try
         {
-            publishResults(project)
+            publishResults(project, skipHostTest)
         }
         catch(ee)
         {}
 
         throw e
     }
-    publishResults(project)
+    publishResults(project, skipHostTest)
 }
 
 return this

.jenkins/integration.groovy

@@ -0,0 +1,82 @@
+#!/usr/bin/env groovy
+// This shared library is available at https://github.com/ROCmSoftwarePlatform/rocJENKINS/
+@Library('rocJenkins@pong') _
+
+// This is file for internal AMD use.
+// If you are interested in running your own Jenkins, please raise a github issue for assistance.
+
+import com.amd.project.*
+import com.amd.docker.*
+import java.nio.file.Path
+
+def runCI =
+{
+    nodeDetails, jobName ->
+
+    def prj = new rocProject('Tensile', 'Integration')
+
+    // Define test architectures, optional rocm version argument is available
+    def nodes = new dockerNodes(nodeDetails, jobName, prj)
+
+    boolean formatCheck = false
+
+    prj.timeout.test = 60
+
+    def commonGroovy
+
+    def compileCommand =
+    {
+        platform, project->
+
+        commonGroovy = load "${project.paths.project_src_prefix}/.jenkins/common.groovy"
+        // commonGroovy.runCompileCommand(platform, project, jobName, false)
+    }
+
+    def testCommand =
+    {
+        platform, project->
+
+        def test_marks = "integration"
+        boolean skipHostTest = true
+        commonGroovy.runTestCommand(platform, project, jobName, test_marks, skipHostTest)
+    }
+
+    buildProject(prj, formatCheck, nodes.dockerArray, compileCommand, testCommand, null)
+
+}
+
+ci: {
+    String urlJobName = auxiliary.getTopJobName(env.BUILD_URL)
+
+    def propertyList = ["compute-rocm-dkms-no-npi-hipclang":[],
+                        "rocm-docker":[]]
+    propertyList = auxiliary.appendPropertyList(propertyList)
+
+    def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([ubuntu18:['gfx900','gfx906','gfx908']]),
+                       "rocm-docker":([ubuntu18:['gfx908']])]
+
+    propertyList.each
+    {
+        jobName, property->
+        if (urlJobName == jobName)
+            properties(auxiliary.addCommonProperties(property))
+    }
+
+    jobNameList.each
+    {
+        jobName, nodeDetails->
+        if (urlJobName == jobName)
+            stage(jobName) {
+                runCI(nodeDetails, jobName)
+            }
+    }
+
+    // For url job names that are outside of the standardJobNameSet i.e. compute-rocm-dkms-no-npi-1901
+    if(!jobNameList.keySet().contains(urlJobName))
+    {
+        properties(auxiliary.addCommonProperties([pipelineTriggers([cron('0 6 * * 6')])]))
+        stage(urlJobName) {
+            runCI([ubuntu18:['any']], urlJobName)
+        }
+    }
+}

HostLibraryTests/CMakeLists.txt

@@ -50,7 +50,7 @@ if(TENSILE_STATIC_ONLY)
 endif()
 
 if(NOT Tensile_FOUND)
-    find_package(Tensile 4.25.0 EXACT REQUIRED ${TENSILE_COMPONENTS} PATHS "${CMAKE_CURRENT_SOURCE_DIR}/../Tensile")
+    find_package(Tensile 4.26.0 EXACT REQUIRED ${TENSILE_COMPONENTS} PATHS "${CMAKE_CURRENT_SOURCE_DIR}/../Tensile")
 endif()
 
 if(NOT TENSILE_DISABLE_CTEST)
@@ -177,8 +177,9 @@ endif()
 
 if(TENSILE_USE_HIP)
 
-    # Make sure the testing libraries get copied to avoid segfaults in test runs.
-    add_dependencies(TensileTests ${HIP_TEST_COPY_TARGET_DEPS})
+    # Make sure the testing libraries get built & copied before the client is built
+    # to avoid segfaults on gtest integration.
+    add_dependencies(TensileTests ${HIP_TEST_LIBRARY_TARGET_DEPS})
 
     target_link_libraries(TensileTests PRIVATE "-Xlinker --whole-archive")
     target_link_libraries(TensileTests PUBLIC ${HIP_TEST_LIBRARIES})

HostLibraryTests/ContractionProblem_test.cpp

@@ -138,6 +138,22 @@ TEST(ContractionProblem, FromOperationIdentifier)
                                                       2.0);
 
     EXPECT_EQ(problem.operationIdentifier(), identifier);
+
+    identifier         = "Contraction_l_ALik_Bjlk_Cijk_Dijk";
+    auto mirrorProblem = ContractionProblem::FromIndexSizes(identifier,
+                                                            sizes,
+                                                            DataType::Float,
+                                                            empty,
+                                                            DataType::Float,
+                                                            empty,
+                                                            DataType::Float,
+                                                            empty,
+                                                            DataType::Float,
+                                                            empty,
+                                                            2.0);
+    EXPECT_TRUE(mirrorProblem.boundIndices()[0].aMirror);
+    EXPECT_FALSE(mirrorProblem.boundIndices()[0].bMirror);
+    EXPECT_EQ(mirrorProblem.operationIdentifier(), identifier);
 }
 
 #if 0

HostLibraryTests/DataTypes_test.cpp

@@ -45,6 +45,7 @@ using InputTypes = ::testing::Types<std::tuple<float>,
                                     std::tuple<Tensile::BFloat16>,
                                     std::tuple<std::complex<float>>,
                                     std::tuple<std::complex<double>>,
+                                    std::tuple<int8_t>,
                                     std::tuple<Tensile::Int8x4>,
                                     std::tuple<int32_t>>;
 
@@ -84,6 +85,7 @@ static_assert(Tensile::TypeInfo<std::complex<float>>::Enum == Tensile::DataType:
 static_assert(Tensile::TypeInfo<std::complex<double>>::Enum == Tensile::DataType::ComplexDouble,
               "ComplexDouble");
 static_assert(Tensile::TypeInfo<Tensile::Half>::Enum == Tensile::DataType::Half, "Half");
+static_assert(Tensile::TypeInfo<int8_t>::Enum == Tensile::DataType::Int8, "Int8");
 static_assert(Tensile::TypeInfo<Tensile::Int8x4>::Enum == Tensile::DataType::Int8x4, "Int8x4");
 static_assert(Tensile::TypeInfo<int32_t>::Enum == Tensile::DataType::Int32, "Int32");
 static_assert(Tensile::TypeInfo<Tensile::BFloat16>::Enum == Tensile::DataType::BFloat16,
@@ -94,6 +96,7 @@ static_assert(Tensile::TypeInfo<double>::Packing == 1, "Double");
 static_assert(Tensile::TypeInfo<std::complex<float>>::Packing == 1, "ComplexFloat");
 static_assert(Tensile::TypeInfo<std::complex<double>>::Packing == 1, "ComplexDouble");
 static_assert(Tensile::TypeInfo<Tensile::Half>::Packing == 1, "Half");
+static_assert(Tensile::TypeInfo<int8_t>::Packing == 1, "Int8");
 static_assert(Tensile::TypeInfo<Tensile::Int8x4>::Packing == 4, "Int8x4");
 static_assert(Tensile::TypeInfo<int32_t>::Packing == 1, "Int32");
 static_assert(Tensile::TypeInfo<Tensile::BFloat16>::Packing == 1, "BFloat16");
@@ -133,6 +136,7 @@ INSTANTIATE_TEST_SUITE_P(DataTypesTest,
                                            Tensile::DataType::ComplexFloat,
                                            Tensile::DataType::ComplexDouble,
                                            Tensile::DataType::Half,
+                                           Tensile::DataType::Int8,
                                            Tensile::DataType::Int8x4,
                                            Tensile::DataType::Int32,
                                            Tensile::DataType::BFloat16));

HostLibraryTests/hip/CMakeLists.txt

@@ -17,11 +17,13 @@ TensileCreateLibraryFiles(
     CODE_OBJECT_VERSION ${CODE_OBJECT_VERSION}
     )
 
+# Tests need library and kernels copied to data folder
 TensileCreateCopyTarget(
     copy_kernels_lite
     "${LITE_ALL_FILES}"
     ${TEST_DATA_DIR}/kernels_lite
     )
+add_dependencies(copy_kernels_lite LITE_LIBRARY_TARGET)
 
 TensileCreateLibraryFiles(
     "${CMAKE_CURRENT_SOURCE_DIR}/../configs/lite_configs_mixed"
@@ -35,11 +37,13 @@ TensileCreateLibraryFiles(
     CODE_OBJECT_VERSION ${CODE_OBJECT_VERSION}
     )
 
+# Tests need library and kernels copied to data folder
 TensileCreateCopyTarget(
     copy_kernels_lite_mixed
     "${LITE_MIXED_ALL_FILES}"
     ${TEST_DATA_DIR}/kernels_lite_mixed
     )
+add_dependencies(copy_kernels_lite_mixed LITE_MIXED_LIBRARY_TARGET)
 
 TensileCreateLibraryFiles(
     "${CMAKE_CURRENT_SOURCE_DIR}/../configs/lite_configs"
@@ -53,28 +57,33 @@ TensileCreateLibraryFiles(
     CODE_OBJECT_VERSION ${CODE_OBJECT_VERSION}
     )
 
+# Tests need library and kernels copied to data folder
 TensileCreateCopyTarget(
     copy_kernels_lite_2
     "${LITE_2_ALL_FILES}"
     ${TEST_DATA_DIR}/kernels_lite_2
     )
+add_dependencies(copy_kernels_lite_2 LITE_2_LIBRARY_TARGET)
 
 TensileCreateLibraryFiles(
     "${CMAKE_CURRENT_SOURCE_DIR}/../configs/tile_aware_selection"
     "${CMAKE_CURRENT_BINARY_DIR}/test_tile_aware_selection"
     TENSILE_ROOT "${TENSILE_SCRIPT_ROOT}"
     EMBED_LIBRARY test_tile_aware_selection
     EMBED_KEY     tile_aware_selection
+    VAR_PREFIX TILE_AWARE
     NO_MERGE_FILES
     COMPILER ${COMPILER}
     CODE_OBJECT_VERSION ${CODE_OBJECT_VERSION}
     )
 
+# Tests need library and kernels copied to data folder
 TensileCreateCopyTarget(
     copy_kernels_tile_aware_selection
-    "${TENSILE_ALL_FILES}"
+    "${TILE_AWARE_ALL_FILES}"
     ${TEST_DATA_DIR}/tile_aware_selection/library
     )
+add_dependencies(copy_kernels_tile_aware_selection TILE_AWARE_LIBRARY_TARGET)
 
 set(test_yaml rocblas_sgemm_asm_single_kernel.yaml)
 file(COPY ${test_yaml} DESTINATION .)
@@ -97,10 +106,10 @@ set(HIP_TEST_LIBRARIES
 
 # Make sure that the TensileTests depends on targets below.
 # Otherwise the TensileTests executable will segfault out.
-set(HIP_TEST_COPY_TARGET_DEPS
+set(HIP_TEST_LIBRARY_TARGET_DEPS
     copy_kernels_lite
-    copy_kernels_lite_2
     copy_kernels_lite_mixed
+    copy_kernels_lite_2
     copy_kernels_tile_aware_selection
     PARENT_SCOPE
 )