From 0831736e70f59af73ede6b5a943d03d208df15ec 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 --- hipamd/docs/markdown/hip_debugging.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/docs/markdown/hip_debugging.md b/hipamd/docs/markdown/hip_debugging.md index e7e058d17a..bf877d894e 100644 --- a/hipamd/docs/markdown/hip_debugging.md +++ b/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. From 8fae5663d279bfe420aa8968a43b2f812b5bda7a 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. --- hipamd/tests/hipify-clang/run_test.bat | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/hipamd/tests/hipify-clang/run_test.bat b/hipamd/tests/hipify-clang/run_test.bat index 6eefb7e46e..b4858db807 100644 --- a/hipamd/tests/hipify-clang/run_test.bat +++ b/hipamd/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 35bd23fb07a1bd48f09749473827de782aa6eda1 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 --- hipamd/tests/hipify-clang/run_test.bat | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/hipify-clang/run_test.bat b/hipamd/tests/hipify-clang/run_test.bat index b4858db807..d8c8d74cf0 100644 --- a/hipamd/tests/hipify-clang/run_test.bat +++ b/hipamd/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 b9fc653b089e0827cc2f95506bd165e67abe228b 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. --- hipamd/hipify-clang/src/HipifyAction.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..7d9ddebf51 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/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 aaa95597689700a0c0117a7e8fb18e4e7987b600 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 --- hipamd/hipify-clang/src/CUDA2HipMap.cpp | 8 ++--- hipamd/hipify-clang/src/HipifyAction.cpp | 35 ++++++++++++++------ hipamd/hipify-clang/src/HipifyAction.h | 5 +-- hipamd/tests/hipify-clang/headers_test_01.cu | 4 ++- hipamd/tests/hipify-clang/headers_test_02.cu | 10 +++--- hipamd/tests/hipify-clang/headers_test_06.cu | 8 +++++ hipamd/tests/hipify-clang/headers_test_07.cu | 8 +++++ hipamd/tests/hipify-clang/headers_test_08.cu | 14 ++++++++ 8 files changed, 71 insertions(+), 21 deletions(-) create mode 100644 hipamd/tests/hipify-clang/headers_test_06.cu create mode 100644 hipamd/tests/hipify-clang/headers_test_07.cu create mode 100644 hipamd/tests/hipify-clang/headers_test_08.cu diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 1f9c6287ed..6bcc0b38c4 100644 --- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipamd/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/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..9ac0eacbd1 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/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()); @@ -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/hipamd/hipify-clang/src/HipifyAction.h b/hipamd/hipify-clang/src/HipifyAction.h index a269a37117..42622c1e01 100644 --- a/hipamd/hipify-clang/src/HipifyAction.h +++ b/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; /** diff --git a/hipamd/tests/hipify-clang/headers_test_01.cu b/hipamd/tests/hipify-clang/headers_test_01.cu index c39ef80d8f..3747c339e8 100644 --- a/hipamd/tests/hipify-clang/headers_test_01.cu +++ b/hipamd/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/hipamd/tests/hipify-clang/headers_test_02.cu b/hipamd/tests/hipify-clang/headers_test_02.cu index 90d412f797..57308efd59 100644 --- a/hipamd/tests/hipify-clang/headers_test_02.cu +++ b/hipamd/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/hipamd/tests/hipify-clang/headers_test_06.cu b/hipamd/tests/hipify-clang/headers_test_06.cu new file mode 100644 index 0000000000..bce73c42df --- /dev/null +++ b/hipamd/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/hipamd/tests/hipify-clang/headers_test_07.cu b/hipamd/tests/hipify-clang/headers_test_07.cu new file mode 100644 index 0000000000..4237e1eb72 --- /dev/null +++ b/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 +#include "cublas_v2.h" +#include "cublas.h" +#include diff --git a/hipamd/tests/hipify-clang/headers_test_08.cu b/hipamd/tests/hipify-clang/headers_test_08.cu new file mode 100644 index 0000000000..ad54871bd8 --- /dev/null +++ b/hipamd/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 4df35f4c778365d72680bf7615807afa92120a14 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 --- hipamd/tests/hipify-clang/cudaRegister.cu | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/hipamd/tests/hipify-clang/cudaRegister.cu b/hipamd/tests/hipify-clang/cudaRegister.cu index 79d21707c2..43b4345337 100644 --- a/hipamd/tests/hipify-clang/cudaRegister.cu +++ b/hipamd/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 85284dd48fcf5311db775e1dc752ce0389bd0446 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 --- hipamd/Jenkinsfile | 3 ++- hipamd/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index 4909e666f0..6142b94d95 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/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/hipamd/include/hip/hcc_detail/code_object_bundle.hpp b/hipamd/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..2bec0017db 100644 --- a/hipamd/include/hip/hcc_detail/code_object_bundle.hpp +++ b/hipamd/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 6f90e4213f10dae5baeb24b6dd7b28e5564aa448 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 --- hipamd/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/code_object_bundle.hpp b/hipamd/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..1b2dfc1c7c 100644 --- a/hipamd/include/hip/hcc_detail/code_object_bundle.hpp +++ b/hipamd/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 2c89a70d3f0ea7992d2774d0142c62102de13af0 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 --- hipamd/Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index 6142b94d95..6e5f7bc8e5 100644 --- a/hipamd/Jenkinsfile +++ b/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} @@ -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 4995146eaff8358e25c2c7673a4a98266bd92a23 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 --- hipamd/packaging/hip_doc.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/hipamd/packaging/hip_doc.txt b/hipamd/packaging/hip_doc.txt index daef7810b4..5987c2c128 100644 --- a/hipamd/packaging/hip_doc.txt +++ b/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