Merge branch 'master' into feature_purge_genco

hipamd/Jenkinsfile

@@ -151,7 +151,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
       // The rm command needs to run as sudo because the test steps below create files owned by root
       sh  """#!/usr/bin/env bash
           set -x
-          sudo rm -rf ${build_dir_rel}
+          rm -rf ${build_dir_rel}
           mkdir -p ${build_dir_rel}
           cd ${build_dir_rel}
           cmake -DCMAKE_BUILD_TYPE=${build_config} -DCMAKE_INSTALL_PREFIX=staging ${optional_configure} ${source_hip_abs}
@@ -160,6 +160,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
     }
 
     // Cap the maximum amount of testing, in case of hangs
+    // Excluding hipPrintfKernel test from automation; variable fails on CI test machines
     timeout(time: 1, unit: 'HOURS')
     {
       stage("${platform} unit testing")
@@ -169,7 +170,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
             cd ${build_dir_rel}
             make install -j\$(nproc)
             make build_tests -i -j\$(nproc)
-            make test
+            ctest -E hipPrintfKernel
           """
         // If unit tests output a junit or xunit file in the future, jenkins can parse that file
         // to display test results on the dashboard
@@ -193,7 +194,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf
         if( platform.toLowerCase( ).startsWith( 'hcc-ctu' ) )
         {
           archiveArtifacts artifacts: "${build_dir_rel}/*.deb", fingerprint: true
-          archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true
+          // archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true
         }
       }
     }

hipamd/docs/markdown/hip_debugging.md

@@ -74,7 +74,7 @@ HIP provides 3 environment variables in the HIP_*_BLOCKING family.  These introd
 - HIP_API_BLOCKING : Forces hipMemcpyAsync and hipMemsetAsync to be host-synchronous, meaning they will wait for the requested operation to complete before returning to the caller.
 
 These options cause HCC to serialize.  Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP.  
-- HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize.
+- HCC_SERIALIZE_KERNEL : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize.
 - HCC_SERIALIZE_COPY    : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.
 
 - HSA_ENABLE_SDMA=0     : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines.  Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine.  This flag is useful to isolate issues with the hardware copy engines.

hipamd/hipify-clang/src/CUDA2HipMap.cpp

@@ -379,12 +379,12 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
     {"vector_types.h",       {"hip/hip_vector_types.h",   CONV_INCLUDE,             API_RUNTIME}},
 
     // CUBLAS includes
-    {"cublas.h",    {"hipblas.h", CONV_INCLUDE, API_BLAS}},
-    {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}},
+    {"cublas.h",             {"hipblas.h",                CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
+    {"cublas_v2.h",          {"hipblas.h",                CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}},
 
     // CURAND includes
-    {"curand.h",             {"hiprand.h",        CONV_INCLUDE, API_RAND}},
-    {"curand_kernel.h",      {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}},
+    {"curand.h",             {"hiprand.h",                CONV_INCLUDE,             API_RAND}},
+    {"curand_kernel.h",      {"hiprand_kernel.h",         CONV_INCLUDE,             API_RAND}},
 
     // HIP includes
     // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h

hipamd/hipify-clang/src/HipifyAction.cpp

@@ -149,23 +149,38 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
     if (!SM.isWrittenInMainFile(hash_loc)) {
         return;
     }
+    if (!firstHeader) {
+        firstHeader = true;
+        firstHeaderLoc = hash_loc;
+    }
 
     const auto found = CUDA_INCLUDE_MAP.find(file_name);
     if (found == CUDA_INCLUDE_MAP.end()) {
-        if (!firstNotMainHeader) {
-            firstNotMainHeader = true;
-            firstNotMainHeaderLoc = hash_loc;
-        }
         return;
     }
 
     // Special-casing to avoid duplication of the hip_runtime include.
     bool secondMainInclude = false;
-    if (found->second.hipName == "hip/hip_runtime.h") {
-        if (insertedRuntimeHeader) {
-            secondMainInclude = true;
+    if (found->second.countType == CONV_INCLUDE_CUDA_MAIN_H) {
+        switch (found->second.countApiType) {
+            case API_DRIVER:
+            case API_RUNTIME:
+                if (insertedRuntimeHeader) {
+                    secondMainInclude = true;
+                    break;
+                }
+                insertedRuntimeHeader = true;
+                break;
+            case API_BLAS:
+                if (insertedBLASHeader) {
+                    secondMainInclude = true;
+                    break;
+                }
+                insertedBLASHeader = true;
+                break;
+            default:
+                break;
         }
-        insertedRuntimeHeader = true;
     }
 
     Statistics::current().incrementCounter(found->second, file_name.str());
@@ -198,8 +213,13 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
 }
 
 void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) {
-    if (pragmaOnce) { return; }
+    if (pragmaOnce) {
+        return;
+    }
     clang::SourceManager& SM = getCompilerInstance().getSourceManager();
+    if (!SM.isWrittenInMainFile(Loc)) {
+        return;
+    }
     clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
     const clang::Token tok = PP.LookAhead(0);
     StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength());
@@ -356,8 +376,8 @@ void HipifyAction::EndSourceFileAction() {
         clang::SourceLocation sl;
         if (pragmaOnce) {
             sl = pragmaOnceLoc;
-        } else if (firstNotMainHeader) {
-            sl = firstNotMainHeaderLoc;
+        } else if (firstHeader) {
+            sl = firstHeaderLoc;
         } else {
             sl = SM.getLocForStartOfFile(SM.getMainFileID());
         }

hipamd/hipify-clang/src/HipifyAction.h

@@ -23,9 +23,10 @@ private:
     // not, we insert it at the top of the file when we finish processing it.
     // This approach means we do the best it's possible to do w.r.t preserving the user's include order.
     bool insertedRuntimeHeader = false;
-    bool firstNotMainHeader = false;
+    bool insertedBLASHeader = false;
+    bool firstHeader = false;
     bool pragmaOnce = false;
-    clang::SourceLocation firstNotMainHeaderLoc;
+    clang::SourceLocation firstHeaderLoc;
     clang::SourceLocation pragmaOnceLoc;
 
     /**

hipamd/include/hip/hcc_detail/code_object_bundle.hpp

@@ -152,9 +152,9 @@ namespace hip_impl
     };
 
     // CREATORS
-    template<typename I>
-    Bundled_code_header::Bundled_code_header(I f, I l) : Bundled_code_header{}
+    template<typename RandomAccessIterator>
+    Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l) : Bundled_code_header{}
     {
         read(f, l, *this);
     }
-} // Namespace hip_impl.
+} // Namespace hip_impl.

hipamd/packaging/hip_doc.txt

@@ -8,12 +8,12 @@ if(DOXYGEN_EXE)
     install(DIRECTORY RuntimeAPI/html DESTINATION docs/docs/RuntimeAPI)
 endif()
 
-find_program(GRIP_EXE grip)
-if(GRIP_EXE)
-    add_custom_target(convert_md_to_html ALL
-                  COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html)
-    install(DIRECTORY md2html/ DESTINATION docs)
-endif()
+#find_program(GRIP_EXE grip)
+#if(GRIP_EXE)
+#    add_custom_target(convert_md_to_html ALL
+#                  COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html)
+#    install(DIRECTORY md2html/ DESTINATION docs)
+#endif()
 
 #############################
 # Packaging steps

hipamd/tests/hipify-clang/cudaRegister.cu

@@ -22,7 +22,12 @@ THE SOFTWARE.
 #include<cuda.h>
 #include<cuda_runtime.h>
 #include<iostream>
-#include<unistd.h>
+#ifdef _WIN32
+#include <windows.h>
+#define sleep(x) Sleep(x)
+#else
+#include <unistd.h>
+#endif
 #include<stdio.h>
 #include<malloc.h>
 
@@ -33,7 +38,7 @@ THE SOFTWARE.
 // CHECK: if(status != hipSuccess) {
 #define check(msg, status){ \
 if(status != cudaSuccess) { \
-	printf("%s failed. \n", #msg); \
+  printf("%s failed. \n", #msg); \
 } \
 }
 

hipamd/tests/hipify-clang/headers_test_01.cu

@@ -1,6 +1,8 @@
 // RUN: %run_test hipify "%s" "%t" %cuda_args
 
 // CHECK: #include <hip/hip_runtime.h>
+// CHECK-NOT: #include <cuda_runtime.h>
+// CHECK: #include <stdio.h>
 #include <cuda.h>
-// CHECK-NOT: #include<cuda_runtime.h>
 #include <cuda_runtime.h>
+#include <stdio.h>

hipamd/tests/hipify-clang/headers_test_02.cu

@@ -1,6 +1,8 @@
 // RUN: %run_test hipify "%s" "%t" %cuda_args
 
-// CHECK: #include <hip/hip_runtime.h>
-#include <cuda_runtime.h>
-// CHECK-NOT: #include<cuda.h>
-#include <cuda.h>
+// CHECK: #include "hip/hip_runtime.h"
+// CHECK-NOT: #include "cuda_runtime.h"
+// CHECK: #include <stdio.h>
+#include "cuda.h"
+#include "cuda_runtime.h"
+#include <stdio.h>

hipamd/tests/hipify-clang/headers_test_06.cu

@@ -0,0 +1,8 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// CHECK: #include <hipblas.h>
+// CHECK-NOT: #include <cublas_v2.h>
+// CHECK: #include <stdio.h>
+#include <cublas.h>
+#include <cublas_v2.h>
+#include <stdio.h>

hipamd/tests/hipify-clang/headers_test_07.cu

@@ -0,0 +1,8 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// CHECK: #include "hipblas.h"
+// CHECK-NOT: #include "cublas.h"
+// CHECK: #include <stdio.h>
+#include "cublas_v2.h"
+#include "cublas.h"
+#include <stdio.h>

hipamd/tests/hipify-clang/headers_test_08.cu

@@ -0,0 +1,14 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// CHECK: #include <hip/hip_runtime.h>
+// CHECK-NOT: #include <cuda_runtime.h>
+// CHECK: #include <iostream>
+// CHECK: #include "hipblas.h"
+// CHECK-NOT: #include "cublas.h"
+// CHECK: #include <stdio.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <iostream>
+#include "cublas_v2.h"
+#include "cublas.h"
+#include <stdio.h>

hipamd/tests/hipify-clang/run_test.bat

@@ -14,5 +14,6 @@ set clang_args=%4%clang_args%
 
 %HIPIFY% -o=%TMP_FILE% %IN_FILE% -- %clang_args%
 if errorlevel 1 (echo      Error: hipify-clang.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)
-%FILE_CHECK% %IN_FILE% -input-file=%TMP_FILE%
+
+findstr /v /r /c:"[ ]*//[ ]*[CHECK*|RUN]" %TMP_FILE% | %FILE_CHECK% %IN_FILE%
 if errorlevel 1 (echo      Error: FileCheck.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)