From b9a2e08707c97cfd13386ed00e40082cd0774aeb Mon Sep 17 00:00:00 2001 From: Peng Date: Sat, 13 Jan 2018 12:33:27 -0600 Subject: [PATCH 01/10] Update hip_debugging.md on typo of chicken bits [ROCm/hip commit: d9572d7fe7bb3ee556162138d46af09a6024e60e] --- projects/hip/docs/markdown/hip_debugging.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/docs/markdown/hip_debugging.md b/projects/hip/docs/markdown/hip_debugging.md index e7e058d17a..bf877d894e 100644 --- a/projects/hip/docs/markdown/hip_debugging.md +++ b/projects/hip/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. From 974f8022b1973f7e515b57a00a707003f44ac472 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 21:43:18 +0300 Subject: [PATCH 02/10] [HIPIFY][tests][win] Fix run_test.bat All checks should not occur in input file for FileCheck. The issue found on CHECK-NOT. Change removes all lit checks in the hipified file based on regexp, and the resulted stdout is fed as stdin for FileCheck. [ROCm/hip commit: 368db8bf6c66333b989d961b26930fb3da542376] --- projects/hip/tests/hipify-clang/run_test.bat | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/hipify-clang/run_test.bat b/projects/hip/tests/hipify-clang/run_test.bat index 6eefb7e46e..b4858db807 100644 --- a/projects/hip/tests/hipify-clang/run_test.bat +++ b/projects/hip/tests/hipify-clang/run_test.bat @@ -1,4 +1,4 @@ -@echo off +rem @echo off setlocal for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i @@ -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%) From 8f478ccf85927c937cd411543ddf03c8974e2d29 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 21:46:27 +0300 Subject: [PATCH 03/10] [HIPIFY][tests][win] Uncomment @echo off [ROCm/hip commit: c528f4f9c5101f4350c7f31dc141f32d3d8cacaf] --- projects/hip/tests/hipify-clang/run_test.bat | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/run_test.bat b/projects/hip/tests/hipify-clang/run_test.bat index b4858db807..d8c8d74cf0 100644 --- a/projects/hip/tests/hipify-clang/run_test.bat +++ b/projects/hip/tests/hipify-clang/run_test.bat @@ -1,4 +1,4 @@ -rem @echo off +@echo off setlocal for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i From fda3893ee31b5453e540917323ce1c1a38561878 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 23:06:55 +0300 Subject: [PATCH 04/10] [HIPIFY][fix] Fix PragmaDirective File location have to be verified, otherwise location of the first found '#pragma once' in any included header even system will be erroneously handled, which might lead to attempt to including hip_runtime.h in it. [ROCm/hip commit: 77f807b597d788d3cc05752175790261d5697c49] --- projects/hip/hipify-clang/src/HipifyAction.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..7d9ddebf51 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -198,8 +198,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()); From 369f3aa9fad9b508f07ec23781c4a539f96e9a6c Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 23:43:36 +0300 Subject: [PATCH 05/10] [HIPIFY][fix] CUDA and cuBLAS main headers correct handling [ROCm/hip commit: 600d5d7c06f622d6f895fe79fe950f8d0453b4a8] --- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 8 ++--- .../hip/hipify-clang/src/HipifyAction.cpp | 35 +++++++++++++------ projects/hip/hipify-clang/src/HipifyAction.h | 5 +-- .../hip/tests/hipify-clang/headers_test_01.cu | 4 ++- .../hip/tests/hipify-clang/headers_test_02.cu | 10 +++--- .../hip/tests/hipify-clang/headers_test_06.cu | 8 +++++ .../hip/tests/hipify-clang/headers_test_07.cu | 8 +++++ .../hip/tests/hipify-clang/headers_test_08.cu | 14 ++++++++ 8 files changed, 71 insertions(+), 21 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/headers_test_06.cu create mode 100644 projects/hip/tests/hipify-clang/headers_test_07.cu create mode 100644 projects/hip/tests/hipify-clang/headers_test_08.cu diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index 1f9c6287ed..6bcc0b38c4 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -379,12 +379,12 @@ const std::map 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 diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..9ac0eacbd1 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/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()); @@ -356,8 +371,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()); } diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h index a269a37117..42622c1e01 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.h +++ b/projects/hip/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; /** diff --git a/projects/hip/tests/hipify-clang/headers_test_01.cu b/projects/hip/tests/hipify-clang/headers_test_01.cu index c39ef80d8f..3747c339e8 100644 --- a/projects/hip/tests/hipify-clang/headers_test_01.cu +++ b/projects/hip/tests/hipify-clang/headers_test_01.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args // CHECK: #include +// CHECK-NOT: #include +// CHECK: #include #include -// CHECK-NOT: #include #include +#include diff --git a/projects/hip/tests/hipify-clang/headers_test_02.cu b/projects/hip/tests/hipify-clang/headers_test_02.cu index 90d412f797..57308efd59 100644 --- a/projects/hip/tests/hipify-clang/headers_test_02.cu +++ b/projects/hip/tests/hipify-clang/headers_test_02.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args -// CHECK: #include -#include -// CHECK-NOT: #include -#include +// CHECK: #include "hip/hip_runtime.h" +// CHECK-NOT: #include "cuda_runtime.h" +// CHECK: #include +#include "cuda.h" +#include "cuda_runtime.h" +#include diff --git a/projects/hip/tests/hipify-clang/headers_test_06.cu b/projects/hip/tests/hipify-clang/headers_test_06.cu new file mode 100644 index 0000000000..bce73c42df --- /dev/null +++ b/projects/hip/tests/hipify-clang/headers_test_06.cu @@ -0,0 +1,8 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +#include +#include +#include diff --git a/projects/hip/tests/hipify-clang/headers_test_07.cu b/projects/hip/tests/hipify-clang/headers_test_07.cu new file mode 100644 index 0000000000..4237e1eb72 --- /dev/null +++ b/projects/hip/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 +#include "cublas_v2.h" +#include "cublas.h" +#include diff --git a/projects/hip/tests/hipify-clang/headers_test_08.cu b/projects/hip/tests/hipify-clang/headers_test_08.cu new file mode 100644 index 0000000000..ad54871bd8 --- /dev/null +++ b/projects/hip/tests/hipify-clang/headers_test_08.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" +// CHECK: #include +#include +#include +#include +#include "cublas_v2.h" +#include "cublas.h" +#include From 5d83330b47cb131bbe426f1081aed92dd598c9dd Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 24 Jan 2018 20:13:23 +0300 Subject: [PATCH 06/10] [HIPIFY][tests][win] Make cudaRegister.cu building on Windows as well [ROCm/hip commit: f9416a0c49737f0ac19f42565ab6e93a5ea178f1] --- projects/hip/tests/hipify-clang/cudaRegister.cu | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/hipify-clang/cudaRegister.cu b/projects/hip/tests/hipify-clang/cudaRegister.cu index 79d21707c2..43b4345337 100644 --- a/projects/hip/tests/hipify-clang/cudaRegister.cu +++ b/projects/hip/tests/hipify-clang/cudaRegister.cu @@ -22,7 +22,12 @@ THE SOFTWARE. #include #include #include -#include +#ifdef _WIN32 +#include +#define sleep(x) Sleep(x) +#else +#include +#endif #include #include @@ -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); \ } \ } From cc15f73add7c6869aa2b81141f44bcf381cc5e08 Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 24 Jan 2018 17:00:57 -0600 Subject: [PATCH 07/10] Fixing rocblas build failure with ::Bundled_code_header constructor Disabling hipPrintfKernel test from CI [ROCm/hip commit: 0497424978de459d817432efbafb73c2df544b46] --- projects/hip/Jenkinsfile | 3 ++- projects/hip/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 4909e666f0..6142b94d95 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -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 diff --git a/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp b/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..2bec0017db 100644 --- a/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp +++ b/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp @@ -150,8 +150,8 @@ namespace hip_impl }; // CREATORS - template - Bundled_code_header::Bundled_code_header(I f, I l) : Bundled_code_header{} + template + Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l) : Bundled_code_header{} { read(f, l, *this); } From 6e53c6d37a1518b5c3c3c2c2b64e6e82b0685a9c Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 25 Jan 2018 16:51:29 +0000 Subject: [PATCH 08/10] use assign rather than insert [ROCm/hip commit: 8acc8365fa9a424e76c5b1d6d34b121c23ed8de9] --- projects/hip/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp b/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..1b2dfc1c7c 100644 --- a/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp +++ b/projects/hip/include/hip/hcc_detail/code_object_bundle.hpp @@ -88,7 +88,7 @@ namespace hip_impl std::copy_n(it, sizeof(y.cbuf), y.cbuf); it += sizeof(y.cbuf); - y.triple.insert(y.triple.cend(), it, it + y.triple_sz); + y.triple.assign(it, it + y.triple_sz); std::copy_n( f + y.offset, y.bundle_sz, std::back_inserter(y.blob)); @@ -155,4 +155,4 @@ namespace hip_impl { read(f, l, *this); } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. From e987fc9e6f8c75ce94af3f3c8c8ea26f30e4e2ed Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Thu, 25 Jan 2018 10:52:56 -0600 Subject: [PATCH 09/10] Remove archiving of RPM We should archive RPM's from proper centos/fedora machines so that we get the proper dependencies right [ROCm/hip commit: 98f3fe3939148599367f24faf27564ffdf926230] --- projects/hip/Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 6142b94d95..6e5f7bc8e5 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/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} @@ -194,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 } } } From 076ba92854734de79f12fffb42697e8ad055c704 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 26 Jan 2018 07:02:49 +0530 Subject: [PATCH 10/10] Disable md2html conversion in hip doc package [ROCm/hip commit: 9860d5de20116d88943ed6512ebd9ea95d87a22b] --- projects/hip/packaging/hip_doc.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/projects/hip/packaging/hip_doc.txt b/projects/hip/packaging/hip_doc.txt index daef7810b4..5987c2c128 100644 --- a/projects/hip/packaging/hip_doc.txt +++ b/projects/hip/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