From 56f7455515753ffd2fad481a810a173eb57ddb1c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 20 Dec 2017 09:36:00 +0530 Subject: [PATCH 01/37] Return pass on single gpu in hipPeerToPeer_simple [ROCm/hip commit: 9f62283bad4b79d3af2b8245cbbd2fff8d3ca83a] --- .../tests/src/p2p/hipPeerToPeer_simple.cpp | 42 +++++++++---------- 1 file changed, 20 insertions(+), 22 deletions(-) diff --git a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp index 32cab371b8..c279658b5a 100644 --- a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp @@ -397,32 +397,30 @@ int main(int argc, char *argv[]) if (gpuCount < 2) { printf("P2P application requires atleast 2 gpu devices\n"); - return 0; - } + } else { + if (p_tests & 0x100) { + testPeerHostToDevice(false/*useAsyncCopy*/); + } + testPeerHostToDevice(true/*useAsyncCopy*/); - if (p_tests & 0x100) { - testPeerHostToDevice(false/*useAsyncCopy*/); - } - testPeerHostToDevice(true/*useAsyncCopy*/); + if (p_tests & 0x1) { + enablePeerFirst(false/*useAsyncCopy*/); + } - if (p_tests & 0x1) { - enablePeerFirst(false/*useAsyncCopy*/); - } + if (p_tests & 0x2) { + allocMemoryFirst(false/*useAsyncCopy*/); + } - if (p_tests & 0x2) { - allocMemoryFirst(false/*useAsyncCopy*/); - } + if (p_tests & 0x4) { + simpleNegative(); + } - if (p_tests & 0x4) { - simpleNegative(); + if (p_tests & 0x8) { + enablePeerFirst(true/*useAsyncCopy*/); + } + if (p_tests & 0x10) { + allocMemoryFirst(true/*useAsyncCopy*/); + } } - - if (p_tests & 0x8) { - enablePeerFirst(true/*useAsyncCopy*/); - } - if (p_tests & 0x10) { - allocMemoryFirst(true/*useAsyncCopy*/); - } - passed(); } From 75dcb768e03dbc000867dd231838c3aa2a0f3c0d Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 20 Dec 2017 12:05:21 +0530 Subject: [PATCH 02/37] [nvccWarnings] Fix -Wno-deprecated-declarations in hip_anyall and hip_ballot [ROCm/hip commit: bd60ef9031814d3a5e60be0c318e3f97fc4343dd] --- projects/hip/tests/src/deviceLib/hip_anyall.cpp | 2 +- projects/hip/tests/src/deviceLib/hip_ballot.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/deviceLib/hip_anyall.cpp b/projects/hip/tests/src/deviceLib/hip_anyall.cpp index 06354383df..9d455d15b4 100644 --- a/projects/hip/tests/src/deviceLib/hip_anyall.cpp +++ b/projects/hip/tests/src/deviceLib/hip_anyall.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations * RUN: %t * HIT_END */ diff --git a/projects/hip/tests/src/deviceLib/hip_ballot.cpp b/projects/hip/tests/src/deviceLib/hip_ballot.cpp index 14b8f314a1..e4d3cc70a6 100644 --- a/projects/hip/tests/src/deviceLib/hip_ballot.cpp +++ b/projects/hip/tests/src/deviceLib/hip_ballot.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations * RUN: %t * HIT_END */ From 7eb3ff163885dbe18e56f5844c0f4cb3246369c9 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 9 Jan 2018 20:20:28 +0300 Subject: [PATCH 03/37] [HIPIFY][tests] Add Windows testing support [ROCm/hip commit: 3d5b2ef5f68f6824c4b98050aa00dcfeb9d3ffc8] --- projects/hip/tests/hipify-clang/lit.cfg | 17 +++++++++++++---- projects/hip/tests/hipify-clang/lit.site.cfg.in | 13 +++++++++++++ projects/hip/tests/hipify-clang/run_test.bat | 17 +++++++++++++++++ 3 files changed, 43 insertions(+), 4 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/run_test.bat diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 76b2ca08dc..4dca50c052 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -44,8 +44,17 @@ if obj_root is not None: path = os.path.pathsep.join((llvm_tools_dir, config.environment['PATH'])) config.environment['PATH'] = path -config.substitutions.append(("hipify", obj_root+"/hipify-clang")) +hipify_path = obj_root +clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'" -# Clang args for CUDA... -config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30 -isystem%s/samples/common/inc" % (config.cuda_root, config.cuda_root))) -config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) +if sys.platform in ['win32']: + run_test_ext = ".bat" + hipify_path += "/" + config.build_type + clang_args += " -isystem'%s'/common/inc -std=c++14" +else: + run_test_ext = ".sh" + clang_args += " -isystem'%s'/samples/common/inc" + +config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root))) +config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"')) +config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"')) diff --git a/projects/hip/tests/hipify-clang/lit.site.cfg.in b/projects/hip/tests/hipify-clang/lit.site.cfg.in index c1f6804d4d..c1095f65a9 100644 --- a/projects/hip/tests/hipify-clang/lit.site.cfg.in +++ b/projects/hip/tests/hipify-clang/lit.site.cfg.in @@ -1,8 +1,21 @@ import sys +import os config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" +if sys.platform in ['win32']: + config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@" + if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND": + config.cuda_samples_root = os.environ.get('NVCUDASAMPLES_ROOT') + if not config.cuda_samples_root or config.cuda_samples_root == "NVCUDASAMPLES_ROOT-NOTFOUND": + lit_config.fatal('No CUDA Samples dir set! Please set CUDA_SDK_ROOT_DIR.') + config.cuda_sdk_root = config.cuda_samples_root + config.build_type = "@CMAKE_BUILD_TYPE@" + if not config.build_type: + config.build_type = "Debug" +else: + config.cuda_sdk_root = config.cuda_root # Support substitution of the tools and libs dirs with user parameters. This is # used when we can't determine the tool dir at configuration time. diff --git a/projects/hip/tests/hipify-clang/run_test.bat b/projects/hip/tests/hipify-clang/run_test.bat new file mode 100644 index 0000000000..5db0cc5043 --- /dev/null +++ b/projects/hip/tests/hipify-clang/run_test.bat @@ -0,0 +1,17 @@ +@echo off + +for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i +if not defined FILE_CHECK (echo Error: FileCheck.exe not found in PATH. && exit /b 1) + +set HIPIFY=%1 +set IN_FILE=%2 +set TMP_FILE=%3 + +set all_args=%* +call set clang_args=%%all_args:*%4=%% +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% +if errorlevel 1 (echo Error: FileCheck.exe failed with exit code: %errorlevel% && exit /b %errorlevel%) From 23ca29818dcb56547a86645e6e092a09aeb2de42 Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 10 Jan 2018 10:51:01 +0530 Subject: [PATCH 04/37] [nvcc] Enable hipGetDeviceAttribute [ROCm/hip commit: e9f94561660dfb6f43f4da28167a91445a9eca5a] --- .../hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 2919939694..0b965a3ae3 100644 --- a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * RUN: %t EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t * HIT_END */ From 98a7e887a07d359957347e00876722d7678655c7 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 10 Jan 2018 21:03:02 +0300 Subject: [PATCH 05/37] [HIPIFY][tests] Add setlocal to batch script [ROCm/hip commit: fdc10f085650425fe2a8a25bf22f218489b8f62b] --- projects/hip/tests/hipify-clang/run_test.bat | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/tests/hipify-clang/run_test.bat b/projects/hip/tests/hipify-clang/run_test.bat index 5db0cc5043..6eefb7e46e 100644 --- a/projects/hip/tests/hipify-clang/run_test.bat +++ b/projects/hip/tests/hipify-clang/run_test.bat @@ -1,4 +1,5 @@ @echo off +setlocal for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i if not defined FILE_CHECK (echo Error: FileCheck.exe not found in PATH. && exit /b 1) From 6f39820a7cd74170e23b1c4beca9003c1ac305c5 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 10 Jan 2018 21:06:06 +0300 Subject: [PATCH 06/37] [HIPIFY][cmake] Exclude socat from Win config [ROCm/hip commit: ebec2903518151e279dfc8e5ff6d1a624a76b889] --- projects/hip/hipify-clang/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 79d181528b..2b725d7188 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -78,7 +78,9 @@ if (HIPIFY_CLANG_TESTS) require_program(lit) require_program(FileCheck) - require_program(socat) + if(NOT WIN32) + require_program(socat) + endif() # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the # value of --cuda-path for the test runs. From 12ae481f17e9bdce1ad7cc159a93eb6f08294bc8 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 12 Jan 2018 17:15:37 +0300 Subject: [PATCH 07/37] [HIPIFY][tests][win] CUDA samples root env. var is changes Env. var NVCUDASAMPLES_ROOT is changed to NVCUDASAMPLESX_Y_ROOT where X - major ver, Y - minor ver. Reason: NVCUDASAMPLES_ROOT contains path to CUDA SDK installed last, while NVCUDASAMPLESX_Y_ROOT contains samples of the same version as of CUDA_TOOLKIT_ROOT_DIR. [ROCm/hip commit: 5df3a9bb3b3c8f6cced6658ff9be495ca2d70c58] --- projects/hip/tests/hipify-clang/lit.site.cfg.in | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/hipify-clang/lit.site.cfg.in b/projects/hip/tests/hipify-clang/lit.site.cfg.in index c1095f65a9..e52db7b026 100644 --- a/projects/hip/tests/hipify-clang/lit.site.cfg.in +++ b/projects/hip/tests/hipify-clang/lit.site.cfg.in @@ -7,8 +7,10 @@ config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" if sys.platform in ['win32']: config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@" if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND": - config.cuda_samples_root = os.environ.get('NVCUDASAMPLES_ROOT') - if not config.cuda_samples_root or config.cuda_samples_root == "NVCUDASAMPLES_ROOT-NOTFOUND": + cuda_version = "@CUDA_VERSION@" + cuda_version = cuda_version.replace('.','_') + config.cuda_samples_root = os.environ.get('NVCUDASAMPLES' + cuda_version + '_ROOT') + if not config.cuda_samples_root: lit_config.fatal('No CUDA Samples dir set! Please set CUDA_SDK_ROOT_DIR.') config.cuda_sdk_root = config.cuda_samples_root config.build_type = "@CMAKE_BUILD_TYPE@" From ef06c379248b3bddb19b1486dc733ee9cf6afbbd Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 12 Jan 2018 17:36:41 +0300 Subject: [PATCH 08/37] [HIPIFY][cmake] Version compatibility checks of CUDA and clang are added [ROCm/hip commit: caf6b5c4f07e702e9713c3209851eb480ca51306] --- projects/hip/hipify-clang/CMakeLists.txt | 23 +++++++++++++++++++++-- 1 file changed, 21 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 79d181528b..5910698e51 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -65,7 +65,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DI install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) - find_package(PythonInterp 2.7 REQUIRED EXACT) + find_package(PythonInterp 2.7 REQUIRED) function (require_program PROGRAM_NAME) find_program(FOUND_${PROGRAM_NAME} ${PROGRAM_NAME}) @@ -78,11 +78,30 @@ if (HIPIFY_CLANG_TESTS) require_program(lit) require_program(FileCheck) - require_program(socat) + if(NOT WIN32) + require_program(socat) + endif() # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the # value of --cuda-path for the test runs. find_package(CUDA REQUIRED) + if ((CUDA_VERSION VERSION_LESS "7.0") OR (LLVM_PACKAGE_VERSION VERSION_LESS "3.8") OR + (CUDA_VERSION VERSION_GREATER "7.5" AND LLVM_PACKAGE_VERSION VERSION_LESS "4.0") OR + (CUDA_VERSION VERSION_GREATER "8.0" AND LLVM_PACKAGE_VERSION VERSION_LESS "6.0") OR + (CUDA_VERSION VERSION_GREATER "9.0" AND LLVM_PACKAGE_VERSION VERSION_LESS "7.0")) + message(SEND_ERROR "CUDA ${CUDA_VERSION} is not supported by clang ${LLVM_PACKAGE_VERSION}.") + if (CUDA_VERSION VERSION_LESS "7.0") + message(STATUS "Please install CUDA 7.0 or higher.") + elseif ((CUDA_VERSION VERSION_EQUAL "7.0") OR (CUDA_VERSION VERSION_EQUAL "7.5")) + message(STATUS "Please install clang 3.8 or higher.") + elseif (CUDA_VERSION VERSION_EQUAL "8.0") + message(STATUS "Please install clang 4.0 or higher.") + elseif (CUDA_VERSION VERSION_EQUAL "9.0") + message(STATUS "Please install clang 6.0 or higher.") + elseif (CUDA_VERSION VERSION_EQUAL "9.1") + message(STATUS "Please install clang 7.0 or higher.") + endif() + endif() configure_file( ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in From b9a2e08707c97cfd13386ed00e40082cd0774aeb Mon Sep 17 00:00:00 2001 From: Peng Date: Sat, 13 Jan 2018 12:33:27 -0600 Subject: [PATCH 09/37] 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 16d792681ac19ddc3fca425f9f50e17f6844bb80 Mon Sep 17 00:00:00 2001 From: emankov Date: Mon, 15 Jan 2018 14:20:37 +0300 Subject: [PATCH 10/37] [HIPIFY][#311][fix] Get rid of socat in run_test.sh [ROCm/hip commit: 6500cce6aaaf691f6bb9c5e1b805b112b687ceb8] --- projects/hip/hipify-clang/CMakeLists.txt | 3 --- projects/hip/tests/hipify-clang/run_test.sh | 11 ----------- 2 files changed, 14 deletions(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 5910698e51..8b3fa7e591 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/hipify-clang/CMakeLists.txt @@ -78,9 +78,6 @@ if (HIPIFY_CLANG_TESTS) require_program(lit) require_program(FileCheck) - if(NOT WIN32) - require_program(socat) - endif() # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the # value of --cuda-path for the test runs. diff --git a/projects/hip/tests/hipify-clang/run_test.sh b/projects/hip/tests/hipify-clang/run_test.sh index 46b2fc066b..418df5dd4d 100755 --- a/projects/hip/tests/hipify-clang/run_test.sh +++ b/projects/hip/tests/hipify-clang/run_test.sh @@ -13,16 +13,5 @@ shift 3 # Remaining args are the ones to forward to clang proper. -# Time for the classic insane little trick for making colour output work. -# A self-deleting shell-script that does the thing we want to do... -TMP_SCRIPT=$(mktemp) -cat << EOF > $TMP_SCRIPT -set -o errexit -set -o xtrace -rm $TMP_SCRIPT $HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE -EOF -chmod a+x $TMP_SCRIPT -# Run the script via socat, spawning a virtual terminal and propagating exit code, and hence failure. -socat -du EXEC:$TMP_SCRIPT,pty,stderr STDOUT From 0e8da085bf61e820e9b856afc7b242dde02b1d85 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 15 Jan 2018 21:05:05 +0300 Subject: [PATCH 11/37] [HIPIFY][fix][#211] Algorithm for explicit insert of hip include directive If in source CUDA file main header (cuda_runtime.h or cuda.h) is not presented, corresponding HIP main header (hip_runtime.h) should be explicitly included in output hipified file. [Algorithm] 1. If #pragma once is presented, HIP main header should be placed just after it; 2. Otherwise if any other (not CUDA main) header is presented, HIP main header should be placed just before it; 3. Otherwise HIP main header should be placed in the beginning of output file. P.S. There might be one more situation when #ifndef #define ... #endif guard for the entire file is presented (make sense for *.h, *.hpp, *.cuh files). In this case HIP main include should be placed just after such #ifdef, or after #pragma once, if it is also presented. This situation will be handled in a separate change. [ROCm/hip commit: 09655a08530824812349103c377b2d4fba995833] --- .../hip/hipify-clang/src/HipifyAction.cpp | 51 ++++++++++++++----- projects/hip/hipify-clang/src/HipifyAction.h | 9 ++++ .../hip/tests/hipify-clang/headers_test_03.cu | 10 ++++ .../hip/tests/hipify-clang/headers_test_04.cu | 12 +++++ .../hip/tests/hipify-clang/headers_test_05.cu | 12 +++++ 5 files changed, 81 insertions(+), 13 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/headers_test_03.cu create mode 100644 projects/hip/tests/hipify-clang/headers_test_04.cu create mode 100644 projects/hip/tests/hipify-clang/headers_test_05.cu diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 8fb318776d..ada0adec5d 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -152,7 +152,10 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const auto found = CUDA_INCLUDE_MAP.find(file_name); if (found == CUDA_INCLUDE_MAP.end()) { - // Not a CUDA include - don't touch it. + if (!firstNotMainHeader) { + firstNotMainHeader = true; + firstNotMainHeaderLoc = hash_loc; + } return; } @@ -160,7 +163,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, bool secondMainInclude = false; if (found->second.hipName == "hip/hip_runtime.h") { if (insertedRuntimeHeader) { - secondMainInclude = true; + secondMainInclude = true; } insertedRuntimeHeader = true; } @@ -178,15 +181,15 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, // Keep the same include type that the user gave. if (!secondMainInclude) { - clang::SmallString<128> includeBuffer; - if (is_angled) { - newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); - } else { - newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); - } + clang::SmallString<128> includeBuffer; + if (is_angled) { + newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); + } else { + newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); + } } else { - // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. - sl = hash_loc; + // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. + sl = hash_loc; } const char *B = SM.getCharacterData(sl); const char *E = SM.getCharacterData(filename_range.getEnd()); @@ -194,6 +197,18 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, insertReplacement(Rep, clang::FullSourceLoc{sl, SM}); } +void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) { + if (pragmaOnce) { return; } + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + const clang::Token tok = PP.LookAhead(0); + clang::LangOptions DefaultLangOptions; + StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); + if (Text == "once") { + pragmaOnce = true; + pragmaOnceLoc = PP.LookAhead(1).getLocation(); + } +} bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) { StringRef refName = "cudaLaunchKernel"; @@ -339,10 +354,16 @@ void HipifyAction::EndSourceFileAction() { // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert // one copy of the hip include into every file. clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - - clang::SourceLocation sl = SM.getLocForStartOfFile(SM.getMainFileID()); + clang::SourceLocation sl; + if (pragmaOnce) { + sl = pragmaOnceLoc; + } else if (firstNotMainHeader) { + sl = firstNotMainHeaderLoc; + } else { + sl = SM.getLocForStartOfFile(SM.getMainFileID()); + } clang::FullSourceLoc fullSL(sl, SM); - ct::Replacement Rep(SM, sl, 0, "#include \n"); + ct::Replacement Rep(SM, sl, 0, "\n#include \n"); insertReplacement(Rep, fullSL); } @@ -367,6 +388,10 @@ public: const clang::Module* imported) override { hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported); } + + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override { + hipifyAction.PragmaDirective(Loc, Introducer); + } }; } diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h index 03d34601f3..a269a37117 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.h +++ b/projects/hip/hipify-clang/src/HipifyAction.h @@ -23,6 +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 pragmaOnce = false; + clang::SourceLocation firstNotMainHeaderLoc; + clang::SourceLocation pragmaOnceLoc; /** * Rewrite a string literal to refer to hip, not CUDA. @@ -57,6 +61,11 @@ public: StringRef relative_path, const clang::Module *imported); + /** + * Called by the preprocessor for each pragma directive during the non-raw lexing pass. + */ + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer); + protected: /** * Add a Replacement for the current file. These will all be applied after executing the FrontendAction. diff --git a/projects/hip/tests/hipify-clang/headers_test_03.cu b/projects/hip/tests/hipify-clang/headers_test_03.cu new file mode 100644 index 0000000000..0223f11119 --- /dev/null +++ b/projects/hip/tests/hipify-clang/headers_test_03.cu @@ -0,0 +1,10 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #pragma once +// CHECK-NEXT: #include +#pragma once +// CHECK-NOT: #include +int main(int argc, char* argv[]) { + return 0; +} + diff --git a/projects/hip/tests/hipify-clang/headers_test_04.cu b/projects/hip/tests/hipify-clang/headers_test_04.cu new file mode 100644 index 0000000000..f10c0eda56 --- /dev/null +++ b/projects/hip/tests/hipify-clang/headers_test_04.cu @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NEXT: #include +// CHECK-NEXT: #include +#include +#include +// CHECK-NOT: #include +int main(int argc, char* argv[]) { + return 0; +} + diff --git a/projects/hip/tests/hipify-clang/headers_test_05.cu b/projects/hip/tests/hipify-clang/headers_test_05.cu new file mode 100644 index 0000000000..c9428b62d5 --- /dev/null +++ b/projects/hip/tests/hipify-clang/headers_test_05.cu @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #pragma once +// CHECK-NEXT: #include +#pragma once +// CHECK-NOT: #include +#include + +int main(int argc, char* argv[]) { + return 0; +} + From 3be321dfb6d6c341a9ed43c002e63ad35dd1ceba Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 10 Jan 2018 11:26:51 -0600 Subject: [PATCH 12/37] Adding dependencies for rocm_agent_enumerator [ROCm/hip commit: 55fc302e3223c6a3c995143b14a2b7faa89307c5] --- projects/hip/Jenkinsfile | 6 +++--- projects/hip/docker/dockerfile-build-ubuntu-16.04 | 3 +++ 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 12c9755103..ce6fbc2e9e 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -371,7 +371,7 @@ parallel hcc_ctu: { String hcc_ver = 'hcc-ctu' String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/clang_tot_upgrade/hcc-lc-ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd' + String inside_args = '--device=/dev/kfd --device=/dev/dri' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) @@ -411,8 +411,8 @@ hcc_1_6: node('docker && rocm') { String hcc_ver = 'hcc-1.6' - String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/roc-1.6.x/hcc-lc-ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd' + String from_image = 'rocm/dev-ubuntu-16.04:latest' + String inside_args = '--device=/dev/kfd --device=/dev/dri' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) diff --git a/projects/hip/docker/dockerfile-build-ubuntu-16.04 b/projects/hip/docker/dockerfile-build-ubuntu-16.04 index 031bf72437..8f655f7c78 100644 --- a/projects/hip/docker/dockerfile-build-ubuntu-16.04 +++ b/projects/hip/docker/dockerfile-build-ubuntu-16.04 @@ -7,6 +7,7 @@ MAINTAINER Kent Knox ARG user_uid # Install Packages +# python and libnuma1 are dependencies of rocm_agent_enumerator RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \ sudo \ build-essential \ @@ -14,6 +15,8 @@ RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-ins git \ libelf-dev \ rpm \ + python \ + libnuma1 \ && \ apt-get clean && \ rm -rf /var/lib/apt/lists/* From d027fd3394e067c3a8bd64c3044d50cd953659c4 Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 10 Jan 2018 13:28:18 -0600 Subject: [PATCH 13/37] adding group-add flag to docker run [ROCm/hip commit: a9d34afc50f189af8fab23e0e56b099449397e35] --- projects/hip/Jenkinsfile | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index ce6fbc2e9e..4909e666f0 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -367,11 +367,11 @@ if( params.hcc_integration_test ) // The following launches 3 builds in parallel: hcc-ctu, hcc-1.6 and cuda parallel hcc_ctu: { - node('docker && rocm') + node('docker && rocm && dkms') { String hcc_ver = 'hcc-ctu' String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/clang_tot_upgrade/hcc-lc-ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd --device=/dev/dri' + String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) @@ -408,7 +408,7 @@ parallel hcc_ctu: }, hcc_1_6: { - node('docker && rocm') + node('docker && rocm && !dkms') { String hcc_ver = 'hcc-1.6' String from_image = 'rocm/dev-ubuntu-16.04:latest' From e09ba44b1634ab0c84f48056f7ec8b80e0f67bbe Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 15:08:08 +0300 Subject: [PATCH 14/37] Update HipifyAction.cpp dead code eliminate [ROCm/hip commit: e54d9f3df04971b5159a89a830726956500c8369] --- projects/hip/hipify-clang/src/HipifyAction.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index ada0adec5d..7cd5b3d402 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -202,7 +202,6 @@ void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntro clang::SourceManager& SM = getCompilerInstance().getSourceManager(); clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); const clang::Token tok = PP.LookAhead(0); - clang::LangOptions DefaultLangOptions; StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); if (Text == "once") { pragmaOnce = true; From 02bd6e89eb3232a833cfcaf41a8dd1be7e09ba4a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 16:40:31 +0300 Subject: [PATCH 15/37] [HIPIFY][tests] Add more suffixes to lit config [ROCm/hip commit: 58faf1ddc9d15f1e447e88eb775661f438ed58f7] --- projects/hip/tests/hipify-clang/lit.cfg | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 4dca50c052..12b1410cee 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -13,7 +13,7 @@ import lit.util config.name = 'hipify' # suffixes: CUDA source is only supported -config.suffixes = ['.cu'] +config.suffixes = ['.cu','.cuh','.cpp','.c','.hpp','.h'] # testFormat: The test format to use to interpret tests. config.test_format = lit.formats.ShTest() From 1c69f54730c243dbc5ff52ba11dea13a0f63e5e0 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 17:13:29 +0300 Subject: [PATCH 16/37] [HIPIFY][tests] Remove checks on cudaBuiltins As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. Other CUDA builtins are not tracked for now. [ROCm/hip commit: 74b489b5b2fcaad1f9f43bd8c3e02ec5b045f938] --- projects/hip/tests/hipify-clang/axpy.cu | 1 - projects/hip/tests/hipify-clang/cudaRegister.cu | 2 -- projects/hip/tests/hipify-clang/square.cu | 2 -- 3 files changed, 5 deletions(-) diff --git a/projects/hip/tests/hipify-clang/axpy.cu b/projects/hip/tests/hipify-clang/axpy.cu index 071c503a35..2e59fc021a 100644 --- a/projects/hip/tests/hipify-clang/axpy.cu +++ b/projects/hip/tests/hipify-clang/axpy.cu @@ -16,7 +16,6 @@ template __global__ void axpy(T a, T *x, T *y) { - // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } diff --git a/projects/hip/tests/hipify-clang/cudaRegister.cu b/projects/hip/tests/hipify-clang/cudaRegister.cu index 80d17f65b9..79d21707c2 100644 --- a/projects/hip/tests/hipify-clang/cudaRegister.cu +++ b/projects/hip/tests/hipify-clang/cudaRegister.cu @@ -38,7 +38,6 @@ if(status != cudaSuccess) { \ } __global__ void Inc1(float *Ad, float *Bd){ - // CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; int tx = threadIdx.x + blockIdx.x * blockDim.x; if(tx < 1 ){ for(int i=0;i __global__ void vector_square(T *C_d, const T *A_d, size_t N) { - // CHECK: size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - // CHECK: size_t stride = hipBlockDim_x * hipGridDim_x; size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; From d5821de893055f7c8d480fdcbf1d70c00225534b Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 17:21:19 +0300 Subject: [PATCH 17/37] [HIPIFY] Add more supported by HIP CUDA RT API Textures and Arrays data types [ROCm/hip commit: eb61038736bea9f9cc992bd63f2119b476a88d65] --- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index 1893f6ce5b..db95c7216b 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -213,6 +213,13 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cudaMipmappedArray_t", {"hipMipmappedArray_t", CONV_MEM, API_RUNTIME}}, {"cudaMipmappedArray_const_t", {"hipMipmappedArray_const_t", CONV_MEM, API_RUNTIME}}, + // defines + {"cudaArrayDefault", {"hipArrayDefault", CONV_MEM, API_RUNTIME}}, + {"cudaArrayLayered", {"hipArrayLayered", CONV_MEM, API_RUNTIME}}, + {"cudaArraySurfaceLoadStore", {"hipArraySurfaceLoadStore", CONV_MEM, API_RUNTIME}}, + {"cudaArrayCubemap", {"hipArrayCubemap", CONV_MEM, API_RUNTIME}}, + {"cudaArrayTextureGather", {"hipArrayTextureGather", CONV_MEM, API_RUNTIME}}, + {"cudaMemoryAdvise", {"hipMemAdvise", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (CUmem_advise) {"cudaMemRangeAttribute", {"hipMemRangeAttribute", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (CUmem_range_attribute) {"cudaMemcpyKind", {"hipMemcpyKind", CONV_MEM, API_RUNTIME}}, @@ -263,6 +270,15 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cudaSurfaceFormatMode", {"hipSurfaceFormatMode", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, + // defines + {"cudaTextureType1D", {"hipTextureType1D", CONV_TEX, API_RUNTIME}}, + {"cudaTextureType2D", {"hipTextureType2D", CONV_TEX, API_RUNTIME}}, + {"cudaTextureType3D", {"hipTextureType3D", CONV_TEX, API_RUNTIME}}, + {"cudaTextureTypeCubemap", {"hipTextureTypeCubemap", CONV_TEX, API_RUNTIME}}, + {"cudaTextureType1DLayered", {"hipTextureType1DLayered", CONV_TEX, API_RUNTIME}}, + {"cudaTextureType2DLayered", {"hipTextureType2DLayered", CONV_TEX, API_RUNTIME}}, + {"cudaTextureTypeCubemapLayered", {"hipTextureTypeCubemapLayered", CONV_TEX, API_RUNTIME}}, + // Inter-Process Communication (IPC) {"cudaIpcEventHandle_t", {"hipIpcEventHandle_t", CONV_TYPE, API_RUNTIME}}, {"cudaIpcEventHandle_st", {"hipIpcEventHandle_t", CONV_TYPE, API_RUNTIME}}, From 3fc406ead3405f3b710d0e301856318ca8dccbbb Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 19:21:14 +0300 Subject: [PATCH 18/37] Update headers_test_04.cu [ROCm/hip commit: a784a9c824563bd1324b8f877e2fab6c56038c06] --- projects/hip/tests/hipify-clang/headers_test_04.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/headers_test_04.cu b/projects/hip/tests/hipify-clang/headers_test_04.cu index f10c0eda56..57667b5a34 100644 --- a/projects/hip/tests/hipify-clang/headers_test_04.cu +++ b/projects/hip/tests/hipify-clang/headers_test_04.cu @@ -5,7 +5,7 @@ // CHECK-NEXT: #include #include #include -// CHECK-NOT: #include +// CHECK-NOT: #include int main(int argc, char* argv[]) { return 0; } From 0b9b9fd98c104d94a98028332fd6ec1f88ea0cfd Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 19:21:59 +0300 Subject: [PATCH 19/37] Update headers_test_03.cu [ROCm/hip commit: c328e3987cd1484821effe62b89a34357a94a4ac] --- projects/hip/tests/hipify-clang/headers_test_03.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/headers_test_03.cu b/projects/hip/tests/hipify-clang/headers_test_03.cu index 0223f11119..5f2e479683 100644 --- a/projects/hip/tests/hipify-clang/headers_test_03.cu +++ b/projects/hip/tests/hipify-clang/headers_test_03.cu @@ -3,7 +3,7 @@ // CHECK: #pragma once // CHECK-NEXT: #include #pragma once -// CHECK-NOT: #include +// CHECK-NOT: #include int main(int argc, char* argv[]) { return 0; } From 944e84b9b7fb30487f6a05134d94632dceca1d42 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 20:41:08 +0300 Subject: [PATCH 20/37] [HIPIFY][tests] remove concurentKernels.cu as it is one of CUDA SDK samples. [ROCm/hip commit: 9c0dc5cfe78b2fba5de5a85a96c391c11203e5c4] --- .../tests/hipify-clang/concurentKernels.cu | 242 ------------------ 1 file changed, 242 deletions(-) delete mode 100644 projects/hip/tests/hipify-clang/concurentKernels.cu diff --git a/projects/hip/tests/hipify-clang/concurentKernels.cu b/projects/hip/tests/hipify-clang/concurentKernels.cu deleted file mode 100644 index 27e9e0d0e7..0000000000 --- a/projects/hip/tests/hipify-clang/concurentKernels.cu +++ /dev/null @@ -1,242 +0,0 @@ -// RUN: %run_test hipify "%s" "%t" %cuda_args -/* - * Copyright 1993-2015 NVIDIA Corporation. All rights reserved. - * - * Please refer to the NVIDIA end user license agreement (EULA) associated - * with this source code for terms and conditions that govern your use of - * this software. Any use, reproduction, disclosure, or distribution of - * this software and related documentation outside the terms of the EULA - * is strictly prohibited. - * - */ - -// -// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to -// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced -// in CUDA 3.2. -// -// Devices of compute capability 1.x will run the kernels one after another -// Devices of compute capability 2.0 or higher can overlap the kernels -// -#include -#include -#include - -// This is a kernel that does no real work but runs at least for a specified number of clocks -__global__ void clock_block(clock_t *d_o, clock_t clock_count) -{ - unsigned int start_clock = (unsigned int) clock(); - - clock_t clock_offset = 0; - - while (clock_offset < clock_count) - { - unsigned int end_clock = (unsigned int) clock(); - - // The code below should work like - // this (thanks to modular arithmetics): - // - // clock_offset = (clock_t) (end_clock > start_clock ? - // end_clock - start_clock : - // end_clock + (0xffffffffu - start_clock)); - // - // Indeed, let m = 2^32 then - // end - start = end + m - start (mod m). - - clock_offset = (clock_t)(end_clock - start_clock); - } - - d_o[0] = clock_offset; -} - - -// Single warp reduction kernel -__global__ void sum(clock_t *d_clocks, int N) -{ - __shared__ clock_t s_clocks[32]; - - clock_t my_sum = 0; - - for (int i = threadIdx.x; i < N; i+= blockDim.x) - { - my_sum += d_clocks[i]; - } - - s_clocks[threadIdx.x] = my_sum; - syncthreads(); - - for (int i=16; i>0; i/=2) - { - if (threadIdx.x < i) - { - s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i]; - } - - syncthreads(); - } - - d_clocks[0] = s_clocks[0]; -} - -int main(int argc, char **argv) -{ - int nkernels = 8; // number of concurrent kernels - int nstreams = nkernels + 1; // use one more stream than concurrent kernel - int nbytes = nkernels * sizeof(clock_t); // number of data bytes - float kernel_time = 10; // time the kernel should run in ms - float elapsed_time; // timing variables - int cuda_device = 0; - - printf("[%s] - Starting...\n", argv[0]); - - // get number of kernels if overridden on the command line - if (checkCmdLineFlag(argc, (const char **)argv, "nkernels")) - { - nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels"); - nstreams = nkernels + 1; - } - - // use command-line specified CUDA device, otherwise use device with highest Gflops/s - cuda_device = findCudaDevice(argc, (const char **)argv); - - // CHECK: hipDeviceProp_t deviceProp; - cudaDeviceProp deviceProp; - // CHECK: checkCudaErrors(hipGetDevice(&cuda_device)); - checkCudaErrors(cudaGetDevice(&cuda_device)); - - // CHECK: checkCudaErrors(hipGetDeviceProperties(&deviceProp, cuda_device)); - checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device)); - - if ((deviceProp.concurrentKernels == 0)) - { - printf("> GPU does not support concurrent kernel execution\n"); - printf(" CUDA kernel runs will be serialized\n"); - } - - printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n", - deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount); - - // allocate host memory - clock_t *a = 0; // pointer to the array data in host memory - // CHECK: checkCudaErrors(hipHostMalloc((void **)&a, nbytes)); - checkCudaErrors(cudaMallocHost((void **)&a, nbytes)); - - // allocate device memory - clock_t *d_a = 0; // pointers to data and init value in the device memory - // CHECK: checkCudaErrors(hipMalloc((void **)&d_a, nbytes)); - checkCudaErrors(cudaMalloc((void **)&d_a, nbytes)); - - // CHECK: hipStream_t *streams = (hipStream_t *) malloc(nstreams * sizeof(hipStream_t)); - // allocate and initialize an array of stream handles - cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t)); - - for (int i = 0; i < nstreams; i++) - { - // CHECK: checkCudaErrors(hipStreamCreate(&(streams[i]))); - checkCudaErrors(cudaStreamCreate(&(streams[i]))); - } - - // CHECK: hipEvent_t start_event, stop_event; - // create CUDA event handles - cudaEvent_t start_event, stop_event; - - // CHECK: checkCudaErrors(hipEventCreate(&start_event)); - // CHECK: checkCudaErrors(hipEventCreate(&stop_event)); - checkCudaErrors(cudaEventCreate(&start_event)); - checkCudaErrors(cudaEventCreate(&stop_event)); - - // the events are used for synchronization only and hence do not need to record timings - // this also makes events not introduce global sync points when recorded which is critical to get overlap - - // CHECK: hipEvent_t *kernelEvent; - // CHECK: kernelEvent = (hipEvent_t *) malloc(nkernels * sizeof(hipEvent_t)); - cudaEvent_t *kernelEvent; - kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t)); - - for (int i = 0; i < nkernels; i++) - { - // CHECK: checkCudaErrors(hipEventCreateWithFlags(&(kernelEvent[i]), hipEventDisableTiming)); - checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming)); - } - - ////////////////////////////////////////////////////////////////////// - // time execution with nkernels streams - clock_t total_clocks = 0; -#if defined(__arm__) || defined(__aarch64__) - // the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks. - clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000)); -#else - clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate); -#endif - - // CHECK: hipEventRecord(start_event, 0); - cudaEventRecord(start_event, 0); - - // queue nkernels in separate streams and record when they are done - for (int i=0; i>>(&d_a[i], time_clocks); - total_clocks += time_clocks; - - // CHECK: checkCudaErrors(hipEventRecord(kernelEvent[i], streams[i])); - checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i])); - - // make the last stream wait for the kernel event to be recorded - // CHECK: checkCudaErrors(hipStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); - checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); - } - - // queue a sum kernel and a copy back to host in the last stream. - // the commands in this stream get dispatched as soon as all the kernel events have been recorded - // CHECK: hipLaunchKernelGGL(sum, dim3(1), dim3(32), 0, streams[nstreams-1], d_a, nkernels); - // CHECK: checkCudaErrors(hipMemcpyAsync(a, d_a, sizeof(clock_t), hipMemcpyDeviceToHost, streams[nstreams-1])); - sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels); - checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1])); - - // at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel - - // in this sample we just wait until the GPU is done - // CHECK: checkCudaErrors(hipEventRecord(stop_event, 0)); - // CHECK: checkCudaErrors(hipEventSynchronize(stop_event)); - // CHECK: checkCudaErrors(hipEventElapsedTime(&elapsed_time, start_event, stop_event)); - checkCudaErrors(cudaEventRecord(stop_event, 0)); - checkCudaErrors(cudaEventSynchronize(stop_event)); - checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event)); - - printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f); - printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f); - printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f); - - bool bTestResult = (a[0] > total_clocks); - - // release resources - for (int i = 0; i < nkernels; i++) - { - // CHECK: hipStreamDestroy(streams[i]); - // CHECK: hipEventDestroy(kernelEvent[i]); - cudaStreamDestroy(streams[i]); - cudaEventDestroy(kernelEvent[i]); - } - - free(streams); - free(kernelEvent); - - // CHECK: hipEventDestroy(start_event); - // CHECK: hipEventDestroy(stop_event); - // CHECK: hipHostFree(a); - // CHECK: hipFree(d_a); - cudaEventDestroy(start_event); - cudaEventDestroy(stop_event); - cudaFreeHost(a); - cudaFree(d_a); - - if (!bTestResult) - { - printf("Test failed!\n"); - exit(EXIT_FAILURE); - } - - printf("Test passed\n"); - exit(EXIT_SUCCESS); -} From c52681edf2e02122e532107060070deb0beff672 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 21:07:50 +0300 Subject: [PATCH 21/37] [HIPIFY] Add more supported by HIP CUDA Driver API Arrays data types and functions [ROCm/hip commit: 478fed74fec0232cdac46ed8cc8aad2dc0baf9e8] --- ...A_Driver_API_functions_supported_by_HIP.md | 24 +++++++++---------- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 24 +++++++++---------- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 9a4700b19c..5045c9d81d 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -5,8 +5,8 @@ | **type** | **CUDA** | **HIP** | |-------------:|---------------------------------------------------------------|------------------------------------------------------------| | struct | `CUDA_ARRAY3D_DESCRIPTOR` | | -| struct | `CUDA_ARRAY_DESCRIPTOR` | | -| struct | `CUDA_MEMCPY2D` | | +| struct | `CUDA_ARRAY_DESCRIPTOR` | `HIP_ARRAY_DESCRIPTOR` | +| struct | `CUDA_MEMCPY2D` | `hip_Memcpy2D` | | struct | `CUDA_MEMCPY3D` | | | struct | `CUDA_MEMCPY3D_PEER` | | | struct | `CUDA_POINTER_ATTRIBUTE_P2P_TOKENS` | | @@ -27,15 +27,15 @@ | 0x03 |*`CU_CUBEMAP_FACE_NEGATIVE_Y`* | | | 0x04 |*`CU_CUBEMAP_FACE_POSITIVE_Z`* | | | 0x05 |*`CU_CUBEMAP_FACE_NEGATIVE_Z`* | | -| enum |***`CUarray_format`*** | | -| 0x01 |*`CU_AD_FORMAT_UNSIGNED_INT8`* | | -| 0x02 |*`CU_AD_FORMAT_UNSIGNED_INT16`* | | -| 0x03 |*`CU_AD_FORMAT_UNSIGNED_INT32`* | | -| 0x08 |*`CU_AD_FORMAT_SIGNED_INT8`* | | -| 0x09 |*`CU_AD_FORMAT_SIGNED_INT16`* | | -| 0x0a |*`CU_AD_FORMAT_SIGNED_INT32`* | | -| 0x10 |*`CU_AD_FORMAT_HALF`* | | -| 0x20 |*`CU_AD_FORMAT_FLOAT`* | | +| enum |***`CUarray_format`*** |***`hipArray_format`*** | +| 0x01 |*`CU_AD_FORMAT_UNSIGNED_INT8`* |*`HIP_AD_FORMAT_UNSIGNED_INT8`* | +| 0x02 |*`CU_AD_FORMAT_UNSIGNED_INT16`* |*`HIP_AD_FORMAT_UNSIGNED_INT16`* | +| 0x03 |*`CU_AD_FORMAT_UNSIGNED_INT32`* |*`HIP_AD_FORMAT_UNSIGNED_INT32`* | +| 0x08 |*`CU_AD_FORMAT_SIGNED_INT8`* |*`HIP_AD_FORMAT_SIGNED_INT8`* | +| 0x09 |*`CU_AD_FORMAT_SIGNED_INT16`* |*`HIP_AD_FORMAT_SIGNED_INT16`* | +| 0x0a |*`CU_AD_FORMAT_SIGNED_INT32`* |*`HIP_AD_FORMAT_SIGNED_INT32`* | +| 0x10 |*`CU_AD_FORMAT_HALF`* |*`HIP_AD_FORMAT_HALF`* | +| 0x20 |*`CU_AD_FORMAT_FLOAT`* |*`HIP_AD_FORMAT_FLOAT`* | | enum |***`CUctx_flags`*** | | | 0x00 |*`CU_CTX_SCHED_AUTO`* | | | 0x01 |*`CU_CTX_SCHED_SPIN`* | | @@ -518,7 +518,7 @@ | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------| -| `cuArray3DCreate` | | +| `cuArray3DCreate` | `hipArray3DCreate` | | `cuArray3DGetDescriptor` | | | `cuArrayCreate` | | | `cuArrayDestroy` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index db95c7216b..4f468e7b47 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -10,8 +10,8 @@ const std::map CUDA_TYPE_NAME_MAP{ ///////////////////////////// CUDA DRIVER API ///////////////////////////// {"CUDA_ARRAY3D_DESCRIPTOR", {"HIP_ARRAY3D_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUDA_ARRAY_DESCRIPTOR", {"HIP_ARRAY_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUDA_MEMCPY2D", {"HIP_MEMCPY2D", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUDA_ARRAY_DESCRIPTOR", {"HIP_ARRAY_DESCRIPTOR", CONV_TYPE, API_DRIVER}}, + {"CUDA_MEMCPY2D", {"hip_Memcpy2D", CONV_TYPE, API_DRIVER}}, {"CUDA_MEMCPY3D", {"HIP_MEMCPY3D", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUDA_MEMCPY3D_PEER", {"HIP_MEMCPY3D_PEER", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUDA_POINTER_ATTRIBUTE_P2P_TOKENS", {"HIP_POINTER_ATTRIBUTE_P2P_TOKENS", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -23,7 +23,7 @@ const std::map CUDA_TYPE_NAME_MAP{ {"CUaddress_mode", {"hipAddress_mode", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUarray_cubemap_face", {"hipArray_cubemap_face", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUarray_format", {"hipArray_format", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUarray_format", {"hipArray_format", CONV_TYPE, API_DRIVER}}, {"CUcomputemode", {"hipComputemode", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaComputeMode) {"CUmem_advise", {"hipMemAdvise", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaComputeMode) {"CUmem_range_attribute", {"hipMemRangeAttribute", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // API_RUNTIME ANALOGUE (cudaMemRangeAttribute) @@ -564,14 +564,14 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CU_CUBEMAP_FACE_NEGATIVE_Z", {"HIP_CUBEMAP_FACE_NEGATIVE_Z", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x05 // CUarray_format enum - {"CU_AD_FORMAT_UNSIGNED_INT8", {"HIP_AD_FORMAT_UNSIGNED_INT8", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 - {"CU_AD_FORMAT_UNSIGNED_INT16", {"HIP_AD_FORMAT_UNSIGNED_INT16", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 - {"CU_AD_FORMAT_UNSIGNED_INT32", {"HIP_AD_FORMAT_UNSIGNED_INT32", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03 - {"CU_AD_FORMAT_SIGNED_INT8", {"HIP_AD_FORMAT_SIGNED_INT8", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x08 - {"CU_AD_FORMAT_SIGNED_INT16", {"HIP_AD_FORMAT_SIGNED_INT16", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x09 - {"CU_AD_FORMAT_SIGNED_INT32", {"HIP_AD_FORMAT_SIGNED_INT32", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0a - {"CU_AD_FORMAT_HALF", {"HIP_AD_FORMAT_HALF", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10 - {"CU_AD_FORMAT_FLOAT", {"HIP_AD_FORMAT_FLOAT", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20 + {"CU_AD_FORMAT_UNSIGNED_INT8", {"HIP_AD_FORMAT_UNSIGNED_INT8", CONV_TYPE, API_DRIVER}}, // 0x01 + {"CU_AD_FORMAT_UNSIGNED_INT16", {"HIP_AD_FORMAT_UNSIGNED_INT16", CONV_TYPE, API_DRIVER}}, // 0x02 + {"CU_AD_FORMAT_UNSIGNED_INT32", {"HIP_AD_FORMAT_UNSIGNED_INT32", CONV_TYPE, API_DRIVER}}, // 0x03 + {"CU_AD_FORMAT_SIGNED_INT8", {"HIP_AD_FORMAT_SIGNED_INT8", CONV_TYPE, API_DRIVER}}, // 0x08 + {"CU_AD_FORMAT_SIGNED_INT16", {"HIP_AD_FORMAT_SIGNED_INT16", CONV_TYPE, API_DRIVER}}, // 0x09 + {"CU_AD_FORMAT_SIGNED_INT32", {"HIP_AD_FORMAT_SIGNED_INT32", CONV_TYPE, API_DRIVER}}, // 0x0a + {"CU_AD_FORMAT_HALF", {"HIP_AD_FORMAT_HALF", CONV_TYPE, API_DRIVER}}, // 0x10 + {"CU_AD_FORMAT_FLOAT", {"HIP_AD_FORMAT_FLOAT", CONV_TYPE, API_DRIVER}}, // 0x20 // CUcomputemode enum {"CU_COMPUTEMODE_DEFAULT", {"hipComputeModeDefault", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 0 // API_RUNTIME ANALOGUE (cudaComputeModeDefault = 0) @@ -1063,7 +1063,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cuStreamBatchMemOp", {"hipStreamBatchMemOp", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}}, // // no API_Runtime ANALOGUE // Memory management - {"cuArray3DCreate", {"hipArray3DCreate", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuArray3DCreate", {"hipArray3DCreate", CONV_MEM, API_DRIVER}}, {"cuArray3DGetDescriptor", {"hipArray3DGetDescriptor", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}}, {"cuArrayCreate", {"hipArrayCreate", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}}, {"cuArrayDestroy", {"hipArrayDestroy", CONV_MEM, API_DRIVER, HIP_UNSUPPORTED}}, From 205362ebd6ea7acb02bedf82af27b4fcdc2012d7 Mon Sep 17 00:00:00 2001 From: Ashish Date: Wed, 17 Jan 2018 14:39:42 -0800 Subject: [PATCH 22/37] Fixed the mapping of cudaPointerGetAttributes -> hipPointerGetAttributes in the hipify-perl [ROCm/hip commit: c0bb1d818fa5c801a7dc995083b183c3017bb44d] --- projects/hip/bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -281,7 +281,7 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g; $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; - $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttribute_t/g; + $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttributes/g; $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g; $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; From e9cc6399cbaba616a692c57b3a3c0c459e35e75a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 09:14:16 +0300 Subject: [PATCH 23/37] [HIPIFY][perl] Typo fix [ROCm/hip commit: 356eaa81b872d4ca3591ebcdcfaefb0647c28d23] --- projects/hip/bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -281,7 +281,7 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g; $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; - $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttribute_t/g; + $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttributes/g; $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g; $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; From dad19d57f432787984562bfbda0dbaea3aa8f9e4 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 12:03:03 +0300 Subject: [PATCH 24/37] [HIPIFY] Add CUDA Driver API Texture Ref support (partial) [ROCm/hip commit: 5788ac5d379feb1ddcd88bcd5d91298235931342] --- ...A_Driver_API_functions_supported_by_HIP.md | 88 +++++++++--------- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 90 +++++++++---------- 2 files changed, 89 insertions(+), 89 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 5045c9d81d..c1c30ae019 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -248,42 +248,42 @@ | 0x01 |*`CU_RESOURCE_TYPE_MIPMAPPED_ARRAY`* | | | 0x02 |*`CU_RESOURCE_TYPE_LINEAR`* | | | 0x03 |*`CU_RESOURCE_TYPE_PITCH2D`* | | -| enum |***`CUresourceViewFormat`*** | | -| 0x00 |*`CU_RES_VIEW_FORMAT_NONE`* | | -| 0x01 |*`CU_RES_VIEW_FORMAT_UINT_1X8`* | | -| 0x02 |*`CU_RES_VIEW_FORMAT_UINT_2X8`* | | -| 0x03 |*`CU_RES_VIEW_FORMAT_UINT_4X8`* | | -| 0x04 |*`CU_RES_VIEW_FORMAT_SINT_1X8`* | | -| 0x05 |*`CU_RES_VIEW_FORMAT_SINT_2X8`* | | -| 0x06 |*`CU_RES_VIEW_FORMAT_SINT_4X8`* | | -| 0x07 |*`CU_RES_VIEW_FORMAT_UINT_1X16`* | | -| 0x08 |*`CU_RES_VIEW_FORMAT_UINT_2X16`* | | -| 0x09 |*`CU_RES_VIEW_FORMAT_UINT_4X16`* | | -| 0x0a |*`CU_RES_VIEW_FORMAT_SINT_1X16`* | | -| 0x0b |*`CU_RES_VIEW_FORMAT_SINT_2X16`* | | -| 0x0c |*`CU_RES_VIEW_FORMAT_SINT_4X16`* | | -| 0x0d |*`CU_RES_VIEW_FORMAT_UINT_1X32`* | | -| 0x0e |*`CU_RES_VIEW_FORMAT_UINT_2X32`* | | -| 0x0f |*`CU_RES_VIEW_FORMAT_UINT_4X32`* | | -| 0x10 |*`CU_RES_VIEW_FORMAT_SINT_1X32`* | | -| 0x11 |*`CU_RES_VIEW_FORMAT_SINT_2X32`* | | -| 0x12 |*`CU_RES_VIEW_FORMAT_SINT_4X32`* | | -| 0x13 |*`CU_RES_VIEW_FORMAT_FLOAT_1X16`* | | -| 0x14 |*`CU_RES_VIEW_FORMAT_FLOAT_2X16`* | | -| 0x15 |*`CU_RES_VIEW_FORMAT_FLOAT_4X16`* | | -| 0x16 |*`CU_RES_VIEW_FORMAT_FLOAT_1X32`* | | -| 0x17 |*`CU_RES_VIEW_FORMAT_FLOAT_2X32`* | | -| 0x18 |*`CU_RES_VIEW_FORMAT_FLOAT_4X32`* | | -| 0x19 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC1`* | | -| 0x1a |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* | | -| 0x1b |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* | | -| 0x1c |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC4`* | | -| 0x1d |*`CU_RES_VIEW_FORMAT_SIGNED_BC4`* | | -| 0x1e |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC5`* | | -| 0x1f |*`CU_RES_VIEW_FORMAT_SIGNED_BC5`* | | -| 0x20 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC6H`* | | -| 0x21 |*`CU_RES_VIEW_FORMAT_SIGNED_BC6H`* | | -| 0x22 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC7`* | | +| enum |***`CUresourceViewFormat`*** |***`hipResourceViewFormat`*** | +| 0x00 |*`CU_RES_VIEW_FORMAT_NONE`* |*`hipResViewFormatNone`* | +| 0x01 |*`CU_RES_VIEW_FORMAT_UINT_1X8`* |*`hipResViewFormatUnsignedChar1`* | +| 0x02 |*`CU_RES_VIEW_FORMAT_UINT_2X8`* |*`hipResViewFormatUnsignedChar2`* | +| 0x03 |*`CU_RES_VIEW_FORMAT_UINT_4X8`* |*`hipResViewFormatUnsignedChar4`* | +| 0x04 |*`CU_RES_VIEW_FORMAT_SINT_1X8`* |*`hipResViewFormatSignedChar1`* | +| 0x05 |*`CU_RES_VIEW_FORMAT_SINT_2X8`* |*`hipResViewFormatSignedChar2`* | +| 0x06 |*`CU_RES_VIEW_FORMAT_SINT_4X8`* |*`hipResViewFormatSignedChar4`* | +| 0x07 |*`CU_RES_VIEW_FORMAT_UINT_1X16`* |*`hipResViewFormatUnsignedShort1`* | +| 0x08 |*`CU_RES_VIEW_FORMAT_UINT_2X16`* |*`hipResViewFormatUnsignedShort2`* | +| 0x09 |*`CU_RES_VIEW_FORMAT_UINT_4X16`* |*`hipResViewFormatUnsignedShort4`* | +| 0x0a |*`CU_RES_VIEW_FORMAT_SINT_1X16`* |*`hipResViewFormatSignedShort1`* | +| 0x0b |*`CU_RES_VIEW_FORMAT_SINT_2X16`* |*`hipResViewFormatSignedShort2`* | +| 0x0c |*`CU_RES_VIEW_FORMAT_SINT_4X16`* |*`hipResViewFormatSignedShort4`* | +| 0x0d |*`CU_RES_VIEW_FORMAT_UINT_1X32`* |*`hipResViewFormatUnsignedInt1`* | +| 0x0e |*`CU_RES_VIEW_FORMAT_UINT_2X32`* |*`hipResViewFormatUnsignedInt2`* | +| 0x0f |*`CU_RES_VIEW_FORMAT_UINT_4X32`* |*`hipResViewFormatUnsignedInt4`* | +| 0x10 |*`CU_RES_VIEW_FORMAT_SINT_1X32`* |*`hipResViewFormatSignedInt1`* | +| 0x11 |*`CU_RES_VIEW_FORMAT_SINT_2X32`* |*`hipResViewFormatSignedInt2`* | +| 0x12 |*`CU_RES_VIEW_FORMAT_SINT_4X32`* |*`hipResViewFormatSignedInt4`* | +| 0x13 |*`CU_RES_VIEW_FORMAT_FLOAT_1X16`* |*`hipResViewFormatHalf1`* | +| 0x14 |*`CU_RES_VIEW_FORMAT_FLOAT_2X16`* |*`hipResViewFormatHalf2`* | +| 0x15 |*`CU_RES_VIEW_FORMAT_FLOAT_4X16`* |*`hipResViewFormatHalf4`* | +| 0x16 |*`CU_RES_VIEW_FORMAT_FLOAT_1X32`* |*`hipResViewFormatFloat1`* | +| 0x17 |*`CU_RES_VIEW_FORMAT_FLOAT_2X32`* |*`hipResViewFormatFloat2`* | +| 0x18 |*`CU_RES_VIEW_FORMAT_FLOAT_4X32`* |*`hipResViewFormatFloat4`* | +| 0x19 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC1`* |*`hipResViewFormatUnsignedBlockCompressed1`* | +| 0x1a |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC2`* |*`hipResViewFormatUnsignedBlockCompressed2`* | +| 0x1b |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* |*`hipResViewFormatUnsignedBlockCompressed3`* | +| 0x1c |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC4`* |*`hipResViewFormatUnsignedBlockCompressed4`* | +| 0x1d |*`CU_RES_VIEW_FORMAT_SIGNED_BC4`* |*`hipResViewFormatSignedBlockCompressed4`* | +| 0x1e |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC5`* |*`hipResViewFormatUnsignedBlockCompressed5`* | +| 0x1f |*`CU_RES_VIEW_FORMAT_SIGNED_BC5`* |*`hipResViewFormatSignedBlockCompressed5`* | +| 0x20 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC6H`* |*`hipResViewFormatUnsignedBlockCompressed6H`* | +| 0x21 |*`CU_RES_VIEW_FORMAT_SIGNED_BC6H`* |*`hipResViewFormatSignedBlockCompressed6H`* | +| 0x22 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC7`* |*`hipResViewFormatUnsignedBlockCompressed7`* | | enum |***`CUresult`*** |***`hipError_t`*** | | 0 |*`CUDA_SUCCESS`* |*`hipSuccess`* | | 1 |*`CUDA_ERROR_INVALID_VALUE`* |*`hipErrorInvalidValue`* | @@ -394,8 +394,8 @@ | typedef | `CUstreamCallback` | `hipStreamCallback_t` | | typedef | `CUsurfObject` | | | typedef | `CUsurfref` | | -| typedef | `CUtexObject` | | -| typedef | `CUtexref` | | +| typedef | `CUtexObject` | `hipTextureObject_t` | +| typedef | `CUtexref` | `textureReference` | | define |`CU_IPC_HANDLE_SIZE` | | | define |`CU_LAUNCH_PARAM_BUFFER_POINTER` | `HIP_LAUNCH_PARAM_BUFFER_POINTER` | | define |`CU_LAUNCH_PARAM_BUFFER_SIZE` | `HIP_LAUNCH_PARAM_BUFFER_SIZE` | @@ -507,7 +507,7 @@ | `cuModuleGetFunction` | `hipModuleGetFunction` | | `cuModuleGetGlobal` | `hipModuleGetGlobal` | | `cuModuleGetSurfRef` | | -| `cuModuleGetTexRef` | | +| `cuModuleGetTexRef` | `hipModuleGetTexRef` | | `cuModuleLoad` | `hipModuleLoad` | | `cuModuleLoadData` | `hipModuleLoadData` | | `cuModuleLoadDataEx` | `hipModuleLoadDataEx` | @@ -674,12 +674,12 @@ | `cuTexRefGetMipmappedArray` | | | `cuTexRefSetAddress` | | | `cuTexRefSetAddress2D` | | -| `cuTexRefSetAddressMode` | | -| `cuTexRefSetArray` | | +| `cuTexRefSetAddressMode` | `hipTexRefSetAddressMode` | +| `cuTexRefSetArray` | `hipTexRefSetArray` | | `cuTexRefSetBorderColor` | | -| `cuTexRefSetFilterMode` | | -| `cuTexRefSetFlags` | | -| `cuTexRefSetFormat` | | +| `cuTexRefSetFilterMode` | `hipTexRefSetFilterMode` | +| `cuTexRefSetFlags` | `hipTexRefSetFlags` | +| `cuTexRefSetFormat` | `hipTexRefSetFormat` | | `cuTexRefSetMaxAnisotropy` | | | `cuTexRefSetMipmapFilterMode` | | | `cuTexRefSetMipmapLevelBias` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..24c97258cf 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -101,8 +101,8 @@ const std::map CUDA_TYPE_NAME_MAP{ {"CUresourcetype", {"hipResourceType", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaResourceType) {"CUresourcetype_enum", {"hipResourceType", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaResourceType) - {"CUresourceViewFormat", {"hipResourceViewFormat", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaResourceViewFormat) - {"CUresourceViewFormat_enum", {"hipResourceViewFormat", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaResourceViewFormat) + {"CUresourceViewFormat", {"hipResourceViewFormat", CONV_TEX, API_DRIVER}}, // API_Runtime ANALOGUE (cudaResourceViewFormat) + {"CUresourceViewFormat_enum", {"hipResourceViewFormat", CONV_TEX, API_DRIVER}}, // API_Runtime ANALOGUE (cudaResourceViewFormat) {"CUsharedconfig", {"hipSharedMemConfig", CONV_TYPE, API_DRIVER}}, {"CUsharedconfig_enum", {"hipSharedMemConfig", CONV_TYPE, API_DRIVER}}, @@ -128,9 +128,9 @@ const std::map CUDA_TYPE_NAME_MAP{ // typedef struct CUsurfref_st *CUsurfref; {"CUsurfref", {"hipSurfaceReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // {"CUsurfref_st", {"ihipSurfaceReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUtexObject", {"hipTextureObject", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUtexObject", {"hipTextureObject_t", CONV_TYPE, API_DRIVER}}, // typedef struct CUtexref_st *CUtexref; - {"CUtexref", {"hipTextureReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUtexref", {"textureReference", CONV_TYPE, API_DRIVER}}, // {"CUtexref_st", {"ihipTextureReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // Stream Flags enum @@ -871,41 +871,41 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CU_RESOURCE_TYPE_PITCH2D", {"hipResourceTypePitch2D", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03 // API_Runtime ANALOGUE (cudaResourceTypePitch2D = 0x03) // enum CUresourceViewFormat/CUresourceViewFormat_enum - {"CU_RES_VIEW_FORMAT_NONE", {"hipResViewFormatNone", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00 // API_Runtime ANALOGUE (cudaResViewFormatNone = 0x00) - {"CU_RES_VIEW_FORMAT_UINT_1X8", {"hipResViewFormatUnsignedChar1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar1 = 0x01) - {"CU_RES_VIEW_FORMAT_UINT_2X8", {"hipResViewFormatUnsignedChar2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar2 = 0x02) - {"CU_RES_VIEW_FORMAT_UINT_4X8", {"hipResViewFormatUnsignedChar4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar4 = 0x03) - {"CU_RES_VIEW_FORMAT_SINT_1X8", {"hipResViewFormatSignedChar1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar1 = 0x04) - {"CU_RES_VIEW_FORMAT_SINT_2X8", {"hipResViewFormatSignedChar2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x05 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar2 = 0x05) - {"CU_RES_VIEW_FORMAT_SINT_4X8", {"hipResViewFormatSignedChar4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x06 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar4 = 0x06) - {"CU_RES_VIEW_FORMAT_UINT_1X16", {"hipResViewFormatUnsignedShort1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x07 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort1 = 0x07) - {"CU_RES_VIEW_FORMAT_UINT_2X16", {"hipResViewFormatUnsignedShort2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x08 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort2 = 0x08) - {"CU_RES_VIEW_FORMAT_UINT_4X16", {"hipResViewFormatUnsignedShort4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x09 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort4 = 0x09) - {"CU_RES_VIEW_FORMAT_SINT_1X16", {"hipResViewFormatSignedShort1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0a // API_Runtime ANALOGUE (cudaResViewFormatSignedShort1 = 0x0a) - {"CU_RES_VIEW_FORMAT_SINT_2X16", {"hipResViewFormatSignedShort2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0b // API_Runtime ANALOGUE (cudaResViewFormatSignedShort2 = 0x0b) - {"CU_RES_VIEW_FORMAT_SINT_4X16", {"hipResViewFormatSignedShort4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0c // API_Runtime ANALOGUE (cudaResViewFormatSignedShort4 = 0x0c) - {"CU_RES_VIEW_FORMAT_UINT_1X32", {"hipResViewFormatUnsignedInt1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0d // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt1 = 0x0d) - {"CU_RES_VIEW_FORMAT_UINT_2X32", {"hipResViewFormatUnsignedInt2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0e // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt2 = 0x0e) - {"CU_RES_VIEW_FORMAT_UINT_4X32", {"hipResViewFormatUnsignedInt4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0f // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt4 = 0x0f) - {"CU_RES_VIEW_FORMAT_SINT_1X32", {"hipResViewFormatSignedInt1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt1 = 0x10) - {"CU_RES_VIEW_FORMAT_SINT_2X32", {"hipResViewFormatSignedInt2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x11 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt2 = 0x11) - {"CU_RES_VIEW_FORMAT_SINT_4X32", {"hipResViewFormatSignedInt4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x12 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt4 = 0x12) - {"CU_RES_VIEW_FORMAT_FLOAT_1X16", {"hipResViewFormatHalf1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x13 // API_Runtime ANALOGUE (cudaResViewFormatHalf1 = 0x13) - {"CU_RES_VIEW_FORMAT_FLOAT_2X16", {"hipResViewFormatHalf2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x14 // API_Runtime ANALOGUE (cudaResViewFormatHalf2 = 0x14) - {"CU_RES_VIEW_FORMAT_FLOAT_4X16", {"hipResViewFormatHalf4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x15 // API_Runtime ANALOGUE (cudaResViewFormatHalf4 = 0x15) - {"CU_RES_VIEW_FORMAT_FLOAT_1X32", {"hipResViewFormatFloat1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x16 // API_Runtime ANALOGUE (cudaResViewFormatFloat1 = 0x16) - {"CU_RES_VIEW_FORMAT_FLOAT_2X32", {"hipResViewFormatFloat2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x17 // API_Runtime ANALOGUE (cudaResViewFormatFloat2 = 0x17) - {"CU_RES_VIEW_FORMAT_FLOAT_4X32", {"hipResViewFormatFloat4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x18 // API_Runtime ANALOGUE (cudaResViewFormatFloat4 = 0x18) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC1", {"hipResViewFormatUnsignedBlockCompressed1", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x19 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed1 = 0x19) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC2", {"hipResViewFormatUnsignedBlockCompressed2", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1a // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed2 = 0x1a) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC3", {"hipResViewFormatUnsignedBlockCompressed3", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1b // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed3 = 0x1b) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC4", {"hipResViewFormatUnsignedBlockCompressed4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1c // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed4 = 0x1c) - {"CU_RES_VIEW_FORMAT_SIGNED_BC4", {"hipResViewFormatSignedBlockCompressed4", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1d // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed4 = 0x1d) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC5", {"hipResViewFormatUnsignedBlockCompressed5", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1e // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed5 = 0x1e) - {"CU_RES_VIEW_FORMAT_SIGNED_BC5", {"hipResViewFormatSignedBlockCompressed5", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1f // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed5 = 0x1f) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC6H", {"hipResViewFormatUnsignedBlockCompressed6H", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed6H = 0x20) - {"CU_RES_VIEW_FORMAT_SIGNED_BC6H", {"hipResViewFormatSignedBlockCompressed6H", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x21 // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed6H = 0x21) - {"CU_RES_VIEW_FORMAT_UNSIGNED_BC7", {"hipResViewFormatUnsignedBlockCompressed7", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // 0x22 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed7 = 0x22) + {"CU_RES_VIEW_FORMAT_NONE", {"hipResViewFormatNone", CONV_TEX, API_DRIVER}}, // 0x00 // API_Runtime ANALOGUE (cudaResViewFormatNone = 0x00) + {"CU_RES_VIEW_FORMAT_UINT_1X8", {"hipResViewFormatUnsignedChar1", CONV_TEX, API_DRIVER}}, // 0x01 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar1 = 0x01) + {"CU_RES_VIEW_FORMAT_UINT_2X8", {"hipResViewFormatUnsignedChar2", CONV_TEX, API_DRIVER}}, // 0x02 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar2 = 0x02) + {"CU_RES_VIEW_FORMAT_UINT_4X8", {"hipResViewFormatUnsignedChar4", CONV_TEX, API_DRIVER}}, // 0x03 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedChar4 = 0x03) + {"CU_RES_VIEW_FORMAT_SINT_1X8", {"hipResViewFormatSignedChar1", CONV_TEX, API_DRIVER}}, // 0x04 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar1 = 0x04) + {"CU_RES_VIEW_FORMAT_SINT_2X8", {"hipResViewFormatSignedChar2", CONV_TEX, API_DRIVER}}, // 0x05 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar2 = 0x05) + {"CU_RES_VIEW_FORMAT_SINT_4X8", {"hipResViewFormatSignedChar4", CONV_TEX, API_DRIVER}}, // 0x06 // API_Runtime ANALOGUE (cudaResViewFormatSignedChar4 = 0x06) + {"CU_RES_VIEW_FORMAT_UINT_1X16", {"hipResViewFormatUnsignedShort1", CONV_TEX, API_DRIVER}}, // 0x07 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort1 = 0x07) + {"CU_RES_VIEW_FORMAT_UINT_2X16", {"hipResViewFormatUnsignedShort2", CONV_TEX, API_DRIVER}}, // 0x08 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort2 = 0x08) + {"CU_RES_VIEW_FORMAT_UINT_4X16", {"hipResViewFormatUnsignedShort4", CONV_TEX, API_DRIVER}}, // 0x09 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedShort4 = 0x09) + {"CU_RES_VIEW_FORMAT_SINT_1X16", {"hipResViewFormatSignedShort1", CONV_TEX, API_DRIVER}}, // 0x0a // API_Runtime ANALOGUE (cudaResViewFormatSignedShort1 = 0x0a) + {"CU_RES_VIEW_FORMAT_SINT_2X16", {"hipResViewFormatSignedShort2", CONV_TEX, API_DRIVER}}, // 0x0b // API_Runtime ANALOGUE (cudaResViewFormatSignedShort2 = 0x0b) + {"CU_RES_VIEW_FORMAT_SINT_4X16", {"hipResViewFormatSignedShort4", CONV_TEX, API_DRIVER}}, // 0x0c // API_Runtime ANALOGUE (cudaResViewFormatSignedShort4 = 0x0c) + {"CU_RES_VIEW_FORMAT_UINT_1X32", {"hipResViewFormatUnsignedInt1", CONV_TEX, API_DRIVER}}, // 0x0d // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt1 = 0x0d) + {"CU_RES_VIEW_FORMAT_UINT_2X32", {"hipResViewFormatUnsignedInt2", CONV_TEX, API_DRIVER}}, // 0x0e // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt2 = 0x0e) + {"CU_RES_VIEW_FORMAT_UINT_4X32", {"hipResViewFormatUnsignedInt4", CONV_TEX, API_DRIVER}}, // 0x0f // API_Runtime ANALOGUE (cudaResViewFormatUnsignedInt4 = 0x0f) + {"CU_RES_VIEW_FORMAT_SINT_1X32", {"hipResViewFormatSignedInt1", CONV_TEX, API_DRIVER}}, // 0x10 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt1 = 0x10) + {"CU_RES_VIEW_FORMAT_SINT_2X32", {"hipResViewFormatSignedInt2", CONV_TEX, API_DRIVER}}, // 0x11 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt2 = 0x11) + {"CU_RES_VIEW_FORMAT_SINT_4X32", {"hipResViewFormatSignedInt4", CONV_TEX, API_DRIVER}}, // 0x12 // API_Runtime ANALOGUE (cudaResViewFormatSignedInt4 = 0x12) + {"CU_RES_VIEW_FORMAT_FLOAT_1X16", {"hipResViewFormatHalf1", CONV_TEX, API_DRIVER}}, // 0x13 // API_Runtime ANALOGUE (cudaResViewFormatHalf1 = 0x13) + {"CU_RES_VIEW_FORMAT_FLOAT_2X16", {"hipResViewFormatHalf2", CONV_TEX, API_DRIVER}}, // 0x14 // API_Runtime ANALOGUE (cudaResViewFormatHalf2 = 0x14) + {"CU_RES_VIEW_FORMAT_FLOAT_4X16", {"hipResViewFormatHalf4", CONV_TEX, API_DRIVER}}, // 0x15 // API_Runtime ANALOGUE (cudaResViewFormatHalf4 = 0x15) + {"CU_RES_VIEW_FORMAT_FLOAT_1X32", {"hipResViewFormatFloat1", CONV_TEX, API_DRIVER}}, // 0x16 // API_Runtime ANALOGUE (cudaResViewFormatFloat1 = 0x16) + {"CU_RES_VIEW_FORMAT_FLOAT_2X32", {"hipResViewFormatFloat2", CONV_TEX, API_DRIVER}}, // 0x17 // API_Runtime ANALOGUE (cudaResViewFormatFloat2 = 0x17) + {"CU_RES_VIEW_FORMAT_FLOAT_4X32", {"hipResViewFormatFloat4", CONV_TEX, API_DRIVER}}, // 0x18 // API_Runtime ANALOGUE (cudaResViewFormatFloat4 = 0x18) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC1", {"hipResViewFormatUnsignedBlockCompressed1", CONV_TEX, API_DRIVER}}, // 0x19 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed1 = 0x19) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC2", {"hipResViewFormatUnsignedBlockCompressed2", CONV_TEX, API_DRIVER}}, // 0x1a // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed2 = 0x1a) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC3", {"hipResViewFormatUnsignedBlockCompressed3", CONV_TEX, API_DRIVER}}, // 0x1b // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed3 = 0x1b) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC4", {"hipResViewFormatUnsignedBlockCompressed4", CONV_TEX, API_DRIVER}}, // 0x1c // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed4 = 0x1c) + {"CU_RES_VIEW_FORMAT_SIGNED_BC4", {"hipResViewFormatSignedBlockCompressed4", CONV_TEX, API_DRIVER}}, // 0x1d // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed4 = 0x1d) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC5", {"hipResViewFormatUnsignedBlockCompressed5", CONV_TEX, API_DRIVER}}, // 0x1e // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed5 = 0x1e) + {"CU_RES_VIEW_FORMAT_SIGNED_BC5", {"hipResViewFormatSignedBlockCompressed5", CONV_TEX, API_DRIVER}}, // 0x1f // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed5 = 0x1f) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC6H", {"hipResViewFormatUnsignedBlockCompressed6H", CONV_TEX, API_DRIVER}}, // 0x20 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed6H = 0x20) + {"CU_RES_VIEW_FORMAT_SIGNED_BC6H", {"hipResViewFormatSignedBlockCompressed6H", CONV_TEX, API_DRIVER}}, // 0x21 // API_Runtime ANALOGUE (cudaResViewFormatSignedBlockCompressed6H = 0x21) + {"CU_RES_VIEW_FORMAT_UNSIGNED_BC7", {"hipResViewFormatUnsignedBlockCompressed7", CONV_TEX, API_DRIVER}}, // 0x22 // API_Runtime ANALOGUE (cudaResViewFormatUnsignedBlockCompressed7 = 0x22) {"CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE", {"hipSharedMemBankSizeDefault", CONV_TYPE, API_DRIVER}}, {"CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE", {"hipSharedMemBankSizeFourByte", CONV_TYPE, API_DRIVER}}, @@ -997,7 +997,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cuModuleGetFunction", {"hipModuleGetFunction", CONV_MODULE, API_DRIVER}}, {"cuModuleGetGlobal_v2", {"hipModuleGetGlobal", CONV_MODULE, API_DRIVER}}, {"cuModuleGetSurfRef", {"hipModuleGetSurfRef", CONV_MODULE, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuModuleGetTexRef", {"hipModuleGetTexRef", CONV_MODULE, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuModuleGetTexRef", {"hipModuleGetTexRef", CONV_MODULE, API_DRIVER}}, {"cuModuleLoad", {"hipModuleLoad", CONV_MODULE, API_DRIVER}}, {"cuModuleLoadData", {"hipModuleLoadData", CONV_MODULE, API_DRIVER}}, {"cuModuleLoadDataEx", {"hipModuleLoadDataEx", CONV_MODULE, API_DRIVER}}, @@ -1155,12 +1155,12 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cuTexRefGetMipmappedArray", {"hipTexRefGetMipmappedArray", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetAddress", {"hipTexRefSetAddress", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetAddress2D", {"hipTexRefSetAddress2D", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefSetAddressMode", {"hipTexRefSetAddressMode", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefSetArray", {"hipTexRefSetArray", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuTexRefSetAddressMode", {"hipTexRefSetAddressMode", CONV_TEX, API_DRIVER}}, + {"cuTexRefSetArray", {"hipTexRefSetArray", CONV_TEX, API_DRIVER}}, {"cuTexRefSetBorderColor", {"hipTexRefSetBorderColor", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, // // no API_Runtime ANALOGUE - {"cuTexRefSetFilterMode", {"hipTexRefSetFilterMode", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefSetFlags", {"hipTexRefSetFlags", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefSetFormat", {"hipTexRefSetFormat", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuTexRefSetFilterMode", {"hipTexRefSetFilterMode", CONV_TEX, API_DRIVER}}, + {"cuTexRefSetFlags", {"hipTexRefSetFlags", CONV_TEX, API_DRIVER}}, + {"cuTexRefSetFormat", {"hipTexRefSetFormat", CONV_TEX, API_DRIVER}}, {"cuTexRefSetMaxAnisotropy", {"hipTexRefSetMaxAnisotropy", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetMipmapFilterMode", {"hipTexRefSetMipmapFilterMode", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetMipmapLevelBias", {"hipTexRefSetMipmapLevelBias", CONV_TEX, API_DRIVER, HIP_UNSUPPORTED}}, From 7b7560f95d846a42c935b74c2667a22b419ddfb4 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 12:28:56 +0300 Subject: [PATCH 25/37] [HIPIFY] Add cudaMalloc3D support [ROCm/hip commit: ff5f964c079030bc1398a274d9751ad242562db4] --- .../markdown/CUDA_Runtime_API_functions_supported_by_HIP.md | 2 +- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index f27f1538b7..84fd0f469d 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -124,7 +124,7 @@ | `cudaHostRegister` | `hipHostRegister` | | `cudaHostUnregister` | `hipHostUnregister` | | `cudaMalloc` | `hipMalloc` | -| `cudaMalloc3D` | | +| `cudaMalloc3D` | `hipMalloc3D` | | `cudaMalloc3DArray` | `hipMalloc3DArray` | | `cudaMallocArray` | `hipMallocArray` | | `cudaMallocHost` | `hipHostMalloc` | diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..b0c4d83706 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -1441,7 +1441,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudaMalloc", {"hipMalloc", CONV_MEM, API_RUNTIME}}, {"cudaMallocHost", {"hipHostMalloc", CONV_MEM, API_RUNTIME}}, {"cudaMallocArray", {"hipMallocArray", CONV_MEM, API_RUNTIME}}, - {"cudaMalloc3D", {"hipMalloc3D", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaMalloc3D", {"hipMalloc3D", CONV_MEM, API_RUNTIME}}, {"cudaMalloc3DArray", {"hipMalloc3DArray", CONV_MEM, API_RUNTIME}}, {"cudaMallocManaged", {"hipMallocManaged", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, {"cudaMallocMipmappedArray", {"hipMallocMipmappedArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, From 4d1fcf52e30e595f05633799fe1cc73256229751 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 19 Jan 2018 17:38:51 +0300 Subject: [PATCH 26/37] [HIPIFY] cuRAND lib support (partial - only Host) [ROCm/hip commit: 8ff99eeadce5b7062531c6eaffbbabe3fab9cd95] --- .../markdown/CURAND_API_supported_by_HIP.md | 132 ++++++++++++++++++ projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 124 +++++++++++++++- projects/hip/hipify-clang/src/Statistics.cpp | 4 +- projects/hip/hipify-clang/src/Statistics.h | 2 + 4 files changed, 259 insertions(+), 3 deletions(-) create mode 100644 projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md diff --git a/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md new file mode 100644 index 0000000000..83c3d4a234 --- /dev/null +++ b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md @@ -0,0 +1,132 @@ +# CURAND API supported by HIP + +## **1. CURAND Data types** + +| **type** | **CUDA** | **HIP** | **HIP value** (if differs) | +|-------------:|---------------------------------------------------------------|------------------------------------------------------------|----------------------------| +| enum |***`curandStatus`*** |***`hiprandStatus`*** | +| enum |***`curandStatus_t`*** |***`hiprandStatus_t`*** | +| 0 |*`CURAND_STATUS_SUCCESS`* |*`HIPRAND_STATUS_SUCCESS`* | +| 100 |*`CURAND_STATUS_VERSION_MISMATCH`* |*`HIPRAND_STATUS_VERSION_MISMATCH`* | +| 101 |*`CURAND_STATUS_NOT_INITIALIZED`* |*`HIPRAND_STATUS_NOT_INITIALIZED`* | +| 102 |*`CURAND_STATUS_ALLOCATION_FAILED`* |*`HIPRAND_STATUS_ALLOCATION_FAILED`* | +| 103 |*`CURAND_STATUS_TYPE_ERROR`* |*`HIPRAND_STATUS_TYPE_ERROR`* | +| 104 |*`CURAND_STATUS_OUT_OF_RANGE`* |*`HIPRAND_STATUS_OUT_OF_RANGE`* | +| 105 |*`CURAND_STATUS_LENGTH_NOT_MULTIPLE`* |*`HIPRAND_STATUS_LENGTH_NOT_MULTIPLE`* | +| 106 |*`CURAND_STATUS_DOUBLE_PRECISION_REQUIRED`* |*`HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED`* | +| 201 |*`CURAND_STATUS_LAUNCH_FAILURE`* |*`HIPRAND_STATUS_LAUNCH_FAILURE`* | +| 202 |*`CURAND_STATUS_PREEXISTING_FAILURE`* |*`HIPRAND_STATUS_PREEXISTING_FAILURE`* | +| 203 |*`CURAND_STATUS_INITIALIZATION_FAILED`* |*`HIPRAND_STATUS_INITIALIZATION_FAILED`* | +| 204 |*`CURAND_STATUS_ARCH_MISMATCH`* |*`HIPRAND_STATUS_ARCH_MISMATCH`* | +| 999 |*`CURAND_STATUS_INTERNAL_ERROR`* |*`HIPRAND_STATUS_INTERNAL_ERROR`* | +| enum |***`curandRngType`*** |***`hiprandRngType`*** | +| enum |***`curandRngType_t`*** |***`hiprandRngType_t`*** | +| 0 |*`CURAND_RNG_TEST`* |*`HIPRAND_RNG_TEST`* | +| 100 |*`CURAND_RNG_PSEUDO_DEFAULT`* |*`HIPRAND_RNG_PSEUDO_DEFAULT`* | 400 | +| 101 |*`CURAND_RNG_PSEUDO_XORWOW`* |*`HIPRAND_RNG_PSEUDO_XORWOW`* | 401 | +| 121 |*`CURAND_RNG_PSEUDO_MRG32K3A`* |*`HIPRAND_RNG_PSEUDO_MRG32K3A`* | 402 | +| 141 |*`CURAND_RNG_PSEUDO_MTGP32`* |*`HIPRAND_RNG_PSEUDO_MTGP32`* | 403 | +| 142 |*`CURAND_RNG_PSEUDO_MT19937`* |*`HIPRAND_RNG_PSEUDO_MT19937`* | 404 | +| 161 |*`CURAND_RNG_PSEUDO_PHILOX4_32_10`* |*`HIPRAND_RNG_PSEUDO_PHILOX4_32_10`* | 405 | +| 200 |*`CURAND_RNG_QUASI_DEFAULT`* |*`HIPRAND_RNG_QUASI_DEFAULT`* | 500 | +| 201 |*`CURAND_RNG_QUASI_SOBOL32`* |*`HIPRAND_RNG_QUASI_SOBOL32`* | 501 | +| 202 |*`CURAND_RNG_QUASI_SCRAMBLED_SOBOL32`* |*`HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32`* | 502 | +| 203 |*`CURAND_RNG_QUASI_SOBOL64`* |*`HIPRAND_RNG_QUASI_SOBOL64`* | 503 | +| 204 |*`CURAND_RNG_QUASI_SCRAMBLED_SOBOL64`* |*`HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64`* | 504 | +| enum |***`curandOrdering`*** | | +| enum |***`curandOrdering_t`*** | | +| 100 |*`CURAND_ORDERING_PSEUDO_BEST`* | | +| 101 |*`CURAND_ORDERING_PSEUDO_DEFAULT`* | | +| 102 |*`CURAND_ORDERING_PSEUDO_SEEDED`* | | +| 201 |*`CURAND_ORDERING_QUASI_DEFAULT`* | | +| enum |***`curandDirectionVectorSet`*** | | +| enum |***`curandDirectionVectorSet_t`*** | | +| 101 |*`CURAND_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 102 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 103 |*`CURAND_DIRECTION_VECTORS_64_JOEKUO6`* | | +| 104 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`* | | +| uint | `curandDirectionVectors32_t` | `hiprandDirectionVectors32_t` | +| uint | `curandDirectionVectors64_t` | | +| struct | `curandGenerator_st` | `hiprandGenerator_st` | +| struct* | `curandGenerator_t` | `hiprandGenerator_t` | +| enum |***`curandOrdering`*** | | +| enum |***`curandOrdering_t`*** | | +| 100 |*`CURAND_ORDERING_PSEUDO_BEST`* | | +| 101 |*`CURAND_ORDERING_PSEUDO_DEFAULT`* | | +| 102 |*`CURAND_ORDERING_PSEUDO_SEEDED`* | | +| 201 |*`CURAND_ORDERING_QUASI_DEFAULT`* | | +| enum |***`curandDirectionVectorSet`*** | | +| enum |***`curandDirectionVectorSet_t`*** | | +| 101 |*`CURAND_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 102 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 103 |*`CURAND_DIRECTION_VECTORS_64_JOEKUO6`* | | +| 104 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`* | | +| uint | `curandDirectionVectors32_t` | `hiprandDirectionVectors32_t` | +| uint | `curandDirectionVectors64_t` | | +| double | `curandDistribution_st` | | +| double | `curandHistogramM2V_st` | | +| double* | `curandDistribution_t` | | +| double* | `curandHistogramM2V_t` | | +| struct | `curandDistributionShift_st` | | +| struct* | `curandDistributionShift_t` | | +| struct | `curandDistributionM2Shift_st` | | +| struct* | `curandDistributionM2Shift_t` | | +| struct | `curandHistogramM2_st` | | +| struct* | `curandHistogramM2_t` | | +| uint | `curandHistogramM2K_st` | | +| uint* | `curandHistogramM2K_t` | | +| struct | `curandDiscreteDistribution_st` | `hiprandDiscreteDistribution_st` | +| struct* | `curandDiscreteDistribution_t` | `hiprandDiscreteDistribution_t` | +| enum |***`curandMethod`*** | | +| enum |***`curandMethod_t`*** | | +| 0 |*`CURAND_CHOOSE_BEST`* | | +| 1 |*`CURAND_ITR`* | | +| 2 |*`CURAND_KNUTH`* | | +| 3 |*`CURAND_HITR`* | | +| 4 |*`CURAND_M1`* | | +| 5 |*`CURAND_M2`* | | +| 6 |*`CURAND_BINARY_SEARCH`* | | +| 7 |*`CURAND_DISCRETE_GAUSS`* | | +| 8 |*`CURAND_REJECTION`* | | +| 9 |*`CURAND_DEVICE_API`* | | +| 10 |*`CURAND_FAST_REJECTION`* | | +| 11 |*`CURAND_3RD`* | | +| 12 |*`CURAND_DEFINITION`* | | +| 13 |*`CURAND_POISSON`* | | + + +## **2. Host API Functions** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|--------------------------------------------| +| `curandCreateGenerator` | `hiprandCreateGenerator` | +| `curandCreateGeneratorHost` | `hiprandCreateGeneratorHost` | +| `curandCreatePoissonDistribution` | `hiprandCreatePoissonDistribution` | +| `curandDestroyDistribution` | `hiprandDestroyDistribution` | +| `curandDestroyGenerator` | `hiprandDestroyGenerator` | +| `curandGenerate` | `hiprandGenerate` | +| `curandGenerateLogNormal` | `hiprandGenerateLogNormal` | +| `curandGenerateLogNormalDouble` | `hiprandGenerateLogNormalDouble` | +| `curandGenerateLongLong` | | +| `curandGenerateNormal` | `hiprandGenerateNormal` | +| `curandGenerateNormalDouble` | `hiprandGenerateNormalDouble` | +| `curandGeneratePoisson` | `hiprandGeneratePoisson` | +| `curandGenerateSeeds` | `hiprandGenerateSeeds` | +| `curandGenerateUniform` | `hiprandGenerateUniform` | +| `curandGenerateUniformDouble` | `hiprandGenerateUniformDouble` | +| `curandGetDirectionVectors32` | | +| `curandGetDirectionVectors64` | | +| `curandGetProperty` | | +| `curandGetScrambleConstants32` | | +| `curandGetScrambleConstants64` | | +| `curandGetVersion` | `hiprandGetVersion` | +| `curandSetGeneratorOffset` | `hiprandSetGeneratorOffset` | +| `curandSetGeneratorOrdering` | | +| `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` | +| `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` | +| `curandSetStream` | `hiprandSetStream` | + +## **3. Device API Functions** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------| diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index 755ce365af..b410ec43f9 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -323,6 +323,36 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cublasPointerMode_t", {"hipblasPointerMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDataType_t", {"hipblasDataType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, + + ///////////////////////////// cuRAND ///////////////////////////// + {"curandStatus", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, + {"curandStatus_t", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, + {"curandRngType", {"hiprandRngType_t", CONV_TYPE, API_RAND}}, + {"curandRngType_t", {"hiprandRngType_t", CONV_TYPE, API_RAND}}, + {"curandGenerator_st", {"hiprandGenerator_st", CONV_TYPE, API_RAND}}, + {"curandGenerator_t", {"hiprandGenerator_t", CONV_TYPE, API_RAND}}, + {"curandDirectionVectorSet", {"hiprandDirectionVectorSet_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDirectionVectorSet_t", {"hiprandDirectionVectorSet_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandOrdering", {"hiprandOrdering_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandOrdering_t", {"hiprandOrdering_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistribution_st", {"hiprandDistribution_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2V_st", {"hiprandDistribution_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistribution_t", {"hiprandDistribution_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2V_t", {"hiprandDistribution_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionShift_st", {"hiprandDistributionShift_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionShift_t", {"hiprandDistributionShift_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionM2Shift_st", {"hiprandDistributionM2Shift_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionM2Shift_t", {"hiprandDistributionM2Shift_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2_st", {"hiprandHistogramM2_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2_t", {"hiprandHistogramM2_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2K_st", {"hiprandHistogramM2K_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2K_t", {"hiprandHistogramM2K_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDiscreteDistribution_st", {"hiprandDiscreteDistribution_st", CONV_TYPE, API_RAND}}, + {"curandDiscreteDistribution_t", {"hiprandDiscreteDistribution_t", CONV_TYPE, API_RAND}}, + {"curandMethod", {"hiprandMethod_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandMethod_t", {"hiprandMethod_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", CONV_TYPE, API_RAND}}, + {"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, }; /// Maps cuda header names to hip header names. @@ -343,6 +373,10 @@ const std::map CUDA_INCLUDE_MAP{ {"cublas.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, + // CURAND includes + {"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 // {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}}, @@ -2660,7 +2694,95 @@ const std::map CUDA_IDENTIFIER_MAP{ // ROTMG {"cublasSrotmg_v2", {"hipblasSrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasDrotmg_v2", {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}} + {"cublasDrotmg_v2", {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + + ///////////////////////////// cuRAND ///////////////////////////// + // RAND function call status types (enum curandStatus) + {"CURAND_STATUS_SUCCESS", {"HIPRAND_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_VERSION_MISMATCH", {"HIPRAND_STATUS_VERSION_MISMATCH", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_NOT_INITIALIZED", {"HIPRAND_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_ALLOCATION_FAILED", {"HIPRAND_STATUS_ALLOCATION_FAILED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_TYPE_ERROR", {"HIPRAND_STATUS_TYPE_ERROR", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_OUT_OF_RANGE", {"HIPRAND_STATUS_OUT_OF_RANGE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_LENGTH_NOT_MULTIPLE", {"HIPRAND_STATUS_LENGTH_NOT_MULTIPLE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_DOUBLE_PRECISION_REQUIRED", {"HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_LAUNCH_FAILURE", {"HIPRAND_STATUS_LAUNCH_FAILURE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_PREEXISTING_FAILURE", {"HIPRAND_STATUS_PREEXISTING_FAILURE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_INITIALIZATION_FAILED", {"HIPRAND_STATUS_INITIALIZATION_FAILED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_ARCH_MISMATCH", {"HIPRAND_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_INTERNAL_ERROR", {"HIPRAND_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_RAND}}, + + // RAND generator types (enum curandRngType) + {"CURAND_RNG_TEST", {"HIPRAND_RNG_TEST", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_DEFAULT", {"HIPRAND_RNG_PSEUDO_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_XORWOW", {"HIPRAND_RNG_PSEUDO_XORWOW", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MRG32K3A", {"HIPRAND_RNG_PSEUDO_MRG32K3A", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MTGP32", {"HIPRAND_RNG_PSEUDO_MTGP32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MT19937", {"HIPRAND_RNG_PSEUDO_MT19937", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_PHILOX4_32_10", {"HIPRAND_RNG_PSEUDO_PHILOX4_32_10", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_DEFAULT", {"HIPRAND_RNG_QUASI_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SOBOL32", {"HIPRAND_RNG_QUASI_SOBOL32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SCRAMBLED_SOBOL32", {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SOBOL64", {"HIPRAND_RNG_QUASI_SOBOL64", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SCRAMBLED_SOBOL64", {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64", CONV_NUMERIC_LITERAL, API_RAND}}, + + // RAND ordering of results in memory (enum curandOrdering) + {"CURAND_ORDERING_PSEUDO_BEST", {"HIPRAND_ORDERING_PSEUDO_BEST", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_PSEUDO_DEFAULT", {"HIPRAND_ORDERING_PSEUDO_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_PSEUDO_SEEDED", {"HIPRAND_ORDERING_PSEUDO_SEEDED", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_QUASI_DEFAULT", {"HIPRAND_ORDERING_QUASI_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND choice of direction vector set (enum curandDirectionVectorSet) + {"CURAND_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND method (enum curandMethod) + {"CURAND_CHOOSE_BEST", {"HIPRAND_CHOOSE_BEST", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ITR", {"HIPRAND_ITR", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_KNUTH", {"HIPRAND_KNUTH", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_HITR", {"HIPRAND_HITR", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_M1", {"HIPRAND_M1", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_M2", {"HIPRAND_M2", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_BINARY_SEARCH", {"HIPRAND_BINARY_SEARCH", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DISCRETE_GAUSS", {"HIPRAND_DISCRETE_GAUSS", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_REJECTION", {"HIPRAND_REJECTION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DEVICE_API", {"HIPRAND_DEVICE_API", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_FAST_REJECTION", {"HIPRAND_FAST_REJECTION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_3RD", {"HIPRAND_3RD", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DEFINITION", {"HIPRAND_DEFINITION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_POISSON", {"HIPRAND_POISSON", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND Host functions + {"curandCreateGenerator", {"hiprandCreateGenerator", CONV_MATH_FUNC, API_RAND}}, + {"curandCreateGeneratorHost", {"hiprandCreateGeneratorHost", CONV_MATH_FUNC, API_RAND}}, + {"curandCreatePoissonDistribution", {"hiprandCreatePoissonDistribution", CONV_MATH_FUNC, API_RAND}}, + {"curandDestroyDistribution", {"hiprandDestroyDistribution", CONV_MATH_FUNC, API_RAND}}, + {"curandDestroyGenerator", {"hiprandDestroyGenerator", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerate", {"hiprandGenerate", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLogNormal", {"hiprandGenerateLogNormal", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLogNormalDouble", {"hiprandGenerateLogNormalDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLongLong", {"hiprandGenerateLongLong", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGenerateNormal", {"hiprandGenerateNormal", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateNormalDouble", {"hiprandGenerateNormalDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGeneratePoisson", {"hiprandGeneratePoisson", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateSeeds", {"hiprandGenerateSeeds", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateUniform", {"hiprandGenerateUniform", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateUniformDouble", {"hiprandGenerateUniformDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGetDirectionVectors32", {"hiprandGetDirectionVectors32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetDirectionVectors64", {"hiprandGetDirectionVectors64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetProperty", {"hiprandGetProperty", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetScrambleConstants32", {"hiprandGetScrambleConstants32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetScrambleConstants64", {"hiprandGetScrambleConstants64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetVersion", {"hiprandGetVersion", CONV_MATH_FUNC, API_RAND}}, + {"curandSetGeneratorOffset", {"hiprandSetGeneratorOffset", CONV_MATH_FUNC, API_RAND}}, + {"curandSetGeneratorOrdering", {"hiprandSetGeneratorOrdering", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", CONV_MATH_FUNC, API_RAND}}, + {"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", CONV_MATH_FUNC, API_RAND}}, + {"curandSetStream", {"hiprandSetStream", CONV_MATH_FUNC, API_RAND}}, + + // RAND Device functions }; const std::map& CUDA_RENAMES_MAP() { diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index 00b8a66f11..e2b3632066 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -5,7 +5,7 @@ const char *counterNames[NUM_CONV_TYPES] = { - "version", "init", "device", "mem", "kern", "coord_func", "math_func", + "version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func", "special_func", "stream", "event", "occupancy", "ctx", "peer", "module", "cache", "exec", "err", "def", "tex", "gl", "graphics", "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", @@ -14,7 +14,7 @@ const char *counterNames[NUM_CONV_TYPES] = { }; const char *apiNames[NUM_API_TYPES] = { - "CUDA Driver API", "CUDA RT API", "CUBLAS API" + "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API" }; namespace { diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index da4e296db0..f160ca7383 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -18,6 +18,7 @@ enum ConvTypes { CONV_KERN, CONV_COORD_FUNC, CONV_MATH_FUNC, + CONV_DEVICE_FUNC, CONV_SPECIAL_FUNC, CONV_STREAM, CONV_EVENT, @@ -54,6 +55,7 @@ enum ApiTypes { API_DRIVER = 0, API_RUNTIME, API_BLAS, + API_RAND, API_LAST }; constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST; From 562844d37809a666677adb3db347114f92c66007 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 19 Jan 2018 21:29:05 +0300 Subject: [PATCH 27/37] [HIPIFY] cuRAND lib support (Device) [ROCm/hip commit: 6e000adde4b6bbd06e17cf9da54ec640d417890a] --- .../markdown/CURAND_API_supported_by_HIP.md | 41 +++++++++++++++++-- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 37 ++++++++++++++++- 2 files changed, 74 insertions(+), 4 deletions(-) diff --git a/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md index 83c3d4a234..900cfadc03 100644 --- a/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md @@ -93,7 +93,14 @@ | 11 |*`CURAND_3RD`* | | | 12 |*`CURAND_DEFINITION`* | | | 13 |*`CURAND_POISSON`* | | - +| struct | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` | +| struct | `curandStateScrambledSobol64_t` | | +| struct | `curandStateSobol64_t` | | +| struct | `curandStateScrambledSobol32_t` | | +| struct | `curandStateSobol32_t` | `hiprandStateSobol32_t` | +| struct | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` | +| struct | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` | +| struct | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` | ## **2. Host API Functions** @@ -128,5 +135,33 @@ ## **3. Device API Functions** -| **CUDA** | **HIP** | -|-----------------------------------------------------------|-------------------------------| +| **CUDA** | **HIP** | +|-----------------------------------------------------------|--------------------------------------------| +| `curand` | `hiprand` | +| `curand_init` | `hiprand_init` | +| `curand_log_normal` | `hiprand_log_normal` | +| `curand_log_normal_double` | `hiprand_log_normal_double` | +| `curand_log_normal2` | `hiprand_log_normal2` | +| `curand_log_normal2_double` | `hiprand_log_normal2_double` | +| `curand_log_normal4` | `hiprand_log_normal4` | +| `curand_log_normal4_double` | `hiprand_log_normal4_double` | +| `curand_mtgp32_single` | | +| `curand_mtgp32_single_specific` | | +| `curand_mtgp32_specific` | | +| `curand_normal` | `hiprand_normal` | +| `curand_normal_double` | `hiprand_normal_double` | +| `curand_normal2` | `hiprand_normal2` | +| `curand_normal2_double` | `hiprand_normal2_double` | +| `curand_normal4` | `hiprand_normal4` | +| `curand_normal4_double` | `hiprand_normal4_double` | +| `curand_uniform2_double` | `hiprand_uniform2_double` | +| `curand_uniform4` | `hiprand_uniform4` | +| `curand_uniform4_double` | `hiprand_uniform4_double` | +| `curand_discrete` | `hiprand_discrete` | +| `curand_discrete4` | `hiprand_discrete4` | +| `curand_poisson` | `hiprand_poisson` | +| `curand_poisson4` | `hiprand_poisson4` | +| `curand_Philox4x32_10` | | +| `skipahead` | `skipahead` | +| `skipahead_sequence` | `skipahead_sequence` | +| `skipahead_subsequence` | `skipahead_subsequence` | diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index b410ec43f9..db78660121 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -324,7 +324,7 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDataType_t", {"hipblasDataType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - ///////////////////////////// cuRAND ///////////////////////////// + ///////////////////////////// cuRAND ///////////////////////////// {"curandStatus", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, {"curandStatus_t", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, {"curandRngType", {"hiprandRngType_t", CONV_TYPE, API_RAND}}, @@ -353,6 +353,15 @@ const std::map CUDA_TYPE_NAME_MAP{ {"curandMethod_t", {"hiprandMethod_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, {"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", CONV_TYPE, API_RAND}}, {"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + // cuRAND types for Device functions + {"curandStateMtgp32_t", {"hiprandStateMtgp32_t", CONV_TYPE, API_RAND}}, + {"curandStateScrambledSobol64_t", {"hiprandStateScrambledSobol64_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandStateSobol64_t", {"hiprandStateSobol64_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandStateScrambledSobol32_t", {"hiprandStateScrambledSobol32_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandStateSobol32_t", {"hiprandStateSobol32_t", CONV_TYPE, API_RAND}}, + {"curandStateMRG32k3a_t", {"hiprandStateMRG32k3a_t", CONV_TYPE, API_RAND}}, + {"curandStatePhilox4_32_10_t", {"hiprandStatePhilox4_32_10_t", CONV_TYPE, API_RAND}}, + {"curandStateXORWOW_t", {"hiprandStateXORWOW_t", CONV_TYPE, API_RAND}}, }; /// Maps cuda header names to hip header names. @@ -2783,6 +2792,32 @@ const std::map CUDA_IDENTIFIER_MAP{ {"curandSetStream", {"hiprandSetStream", CONV_MATH_FUNC, API_RAND}}, // RAND Device functions + {"curand", {"hiprand", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_init", {"hiprand_init", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal", {"hiprand_log_normal", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal_double", {"hiprand_log_normal_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal2", {"hiprand_log_normal2", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal2_double", {"hiprand_log_normal2_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal4", {"hiprand_log_normal4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_log_normal4_double", {"hiprand_log_normal4_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_mtgp32_single", {"hiprand_mtgp32_single", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curand_mtgp32_single_specific", {"hiprand_mtgp32_single_specific", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curand_mtgp32_specific", {"hiprand_mtgp32_specific", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curand_normal", {"hiprand_normal", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_normal_double", {"hiprand_normal_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_normal2", {"hiprand_normal2", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_normal2_double", {"hiprand_normal2_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_normal4", {"hiprand_normal4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_normal4_double", {"hiprand_normal4_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_uniform2_double", {"hiprand_uniform2_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_uniform4", {"hiprand_uniform4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_uniform4_double", {"hiprand_uniform4_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_discrete", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_discrete4", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_poisson", {"hiprand_poisson", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_poisson4", {"hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_Philox4x32_10", {"hiprand_Philox4x32_10", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}}, + // unchanged function names: skipahead, skipahead_sequence, skipahead_subsequence }; const std::map& CUDA_RENAMES_MAP() { From 13aa621ce4944c8197789ab36991fa6c2f3a26dc Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 22 Jan 2018 17:12:02 +0300 Subject: [PATCH 28/37] [HIPIFY] Sync with hipBLAS ToT and CUDA cuBLAS 9.1 [ROCm/hip commit: ddfb1100808591933e66fe24d41098e00a64fe9d] --- projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 97 +++++++++++++------ 1 file changed, 70 insertions(+), 27 deletions(-) diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index db78660121..1f9c6287ed 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -317,10 +317,10 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cublasOperation_t", {"hipblasOperation_t", CONV_TYPE, API_BLAS}}, {"cublasStatus_t", {"hipblasStatus_t", CONV_TYPE, API_BLAS}}, - {"cublasFillMode_t", {"hipblasFillMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasDiagType_t", {"hipblasDiagType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasSideMode_t", {"hipblasSideMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasPointerMode_t", {"hipblasPointerMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasFillMode_t", {"hipblasFillMode_t", CONV_TYPE, API_BLAS}}, + {"cublasDiagType_t", {"hipblasDiagType_t", CONV_TYPE, API_BLAS}}, + {"cublasSideMode_t", {"hipblasSideMode_t", CONV_TYPE, API_BLAS}}, + {"cublasPointerMode_t", {"hipblasPointerMode_t", CONV_TYPE, API_BLAS}}, {"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDataType_t", {"hipblasDataType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, @@ -2063,7 +2063,11 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasSetKernelStream", {"hipblasSetKernelStream", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasGetAtomicsMode", {"hipblasGetAtomicsMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasSetAtomicsMode", {"hipblasSetAtomicsMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - + {"cublasGetMathMode", {"hipblasGetMathMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasSetMathMode", {"hipblasSetMathMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + + + // Blas operations (cublasOperation_t) {"CUBLAS_OP_N", {"HIPBLAS_OP_N", CONV_NUMERIC_LITERAL, API_BLAS}}, {"CUBLAS_OP_T", {"HIPBLAS_OP_T", CONV_NUMERIC_LITERAL, API_BLAS}}, @@ -2080,20 +2084,20 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUBLAS_STATUS_NOT_SUPPORTED", {"HIPBLAS_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Fill Modes (cublasFillMode_t) - {"CUBLAS_FILL_MODE_LOWER", {"HIPBLAS_FILL_MODE_LOWER", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_FILL_MODE_UPPER", {"HIPBLAS_FILL_MODE_UPPER", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, + {"CUBLAS_FILL_MODE_LOWER", {"HIPBLAS_FILL_MODE_LOWER", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_FILL_MODE_UPPER", {"HIPBLAS_FILL_MODE_UPPER", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Diag Types (cublasDiagType_t) - {"CUBLAS_DIAG_NON_UNIT", {"HIPBLAS_DIAG_NON_UNIT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_DIAG_UNIT", {"HIPBLAS_DIAG_UNIT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, + {"CUBLAS_DIAG_NON_UNIT", {"HIPBLAS_DIAG_NON_UNIT", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_DIAG_UNIT", {"HIPBLAS_DIAG_UNIT", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Side Modes (cublasSideMode_t - {"CUBLAS_SIDE_LEFT", {"HIPBLAS_SIDE_LEFT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_SIDE_RIGHT", {"HIPBLAS_SIDE_RIGHT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, + {"CUBLAS_SIDE_LEFT", {"HIPBLAS_SIDE_LEFT", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_SIDE_RIGHT", {"HIPBLAS_SIDE_RIGHT", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Pointer Modes (cublasPointerMode_t) - {"CUBLAS_POINTER_MODE_HOST", {"HIPBLAS_POINTER_MODE_HOST", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, - {"CUBLAS_POINTER_MODE_DEVICE", {"HIPBLAS_POINTER_MODE_DEVICE", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, + {"CUBLAS_POINTER_MODE_HOST", {"HIPBLAS_POINTER_MODE_HOST", CONV_NUMERIC_LITERAL, API_BLAS}}, + {"CUBLAS_POINTER_MODE_DEVICE", {"HIPBLAS_POINTER_MODE_DEVICE", CONV_NUMERIC_LITERAL, API_BLAS}}, // Blas Atomics Modes (cublasAtomicsMode_t) {"CUBLAS_ATOMICS_NOT_ALLOWED", {"HIPBLAS_ATOMICS_NOT_ALLOWED", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED}}, @@ -2111,18 +2115,24 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasSetVector", {"hipblasSetVector", CONV_MATH_FUNC, API_BLAS}}, {"cublasGetVector", {"hipblasGetVector", CONV_MATH_FUNC, API_BLAS}}, + {"cublasSetVectorAsync", {"hipblasSetVectorAsync", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasGetVectorAsync", {"hipblasGetVectorAsync", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasSetMatrix", {"hipblasSetMatrix", CONV_MATH_FUNC, API_BLAS}}, {"cublasGetMatrix", {"hipblasGetMatrix", CONV_MATH_FUNC, API_BLAS}}, - {"cublasGetMatrixAsync", {"hipblasGetMatrixAsync", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasSetMatrixAsync", {"hipblasSetMatrixAsync", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasXerbla", {"hipblasXerbla", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // NRM2 {"cublasSnrm2", {"hipblasSnrm2", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDnrm2", {"hipblasDnrm2", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasScnrm2", {"hipblasScnrm2", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDznrm2", {"hipblasDznrm2", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasNrm2Ex", {"hipblasNrm2Ex", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // DOT {"cublasSdot", {"hipblasSdot", CONV_MATH_FUNC, API_BLAS}}, // there is no such a function in CUDA @@ -2171,8 +2181,8 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasZswap", {"hipblasZswap", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // AMAX - {"cublasIsamax", {"hipblasIsamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasIdamax", {"hipblasIdamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasIsamax", {"hipblasIsamax", CONV_MATH_FUNC, API_BLAS}}, + {"cublasIdamax", {"hipblasIdamax", CONV_MATH_FUNC, API_BLAS}}, {"cublasIcamax", {"hipblasIcamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIzamax", {"hipblasIzamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2217,7 +2227,7 @@ const std::map CUDA_IDENTIFIER_MAP{ // GEMV {"cublasSgemv", {"hipblasSgemv", CONV_MATH_FUNC, API_BLAS}}, // there is no such a function in CUDA - {"cublasSgemvBatched", {"hipblasSgemvBatched", CONV_MATH_FUNC, API_BLAS}}, + {"cublasSgemvBatched", {"hipblasSgemvBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDgemv", {"hipblasDgemv", CONV_MATH_FUNC, API_BLAS}}, {"cublasCgemv", {"hipblasCgemv", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZgemv", {"hipblasZgemv", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2323,14 +2333,25 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasCgemm", {"hipblasCgemm", CONV_MATH_FUNC, API_BLAS}}, {"cublasZgemm", {"hipblasZgemm", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasHgemm", {"hipblasHgemm", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // BATCH GEMM {"cublasSgemmBatched", {"hipblasSgemmBatched", CONV_MATH_FUNC, API_BLAS}}, {"cublasDgemmBatched", {"hipblasDgemmBatched", CONV_MATH_FUNC, API_BLAS}}, + {"cublasHgemmBatched", {"hipblasHgemmBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasCgemmBatched", {"hipblasCgemmBatched", CONV_MATH_FUNC, API_BLAS}}, + {"cublasSgemmStridedBatched", {"hipblasSgemmStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasDgemmStridedBatched", {"hipblasDgemmStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + + {"cublasCgemmBatched", {"hipblasCgemmBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasCgemm3mBatched", {"hipblasCgemm3mBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZgemmBatched", {"hipblasZgemmBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasCgemmStridedBatched", {"hipblasCgemmStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasCgemm3mStridedBatched", {"hipblasCgemm3mStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasZgemmStridedBatched", {"hipblasZgemmStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasHgemmStridedBatched", {"hipblasHgemmStridedBatched", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // SYRK {"cublasSsyrk", {"hipblasSsyrk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDsyrk", {"hipblasDsyrk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2353,7 +2374,6 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasCsyrkx", {"hipblasCsyrkx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZsyrkx", {"hipblasZsyrkx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - // HER2K {"cublasCher2k", {"hipblasCher2k", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZher2k", {"hipblasZher2k", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2392,8 +2412,8 @@ const std::map CUDA_IDENTIFIER_MAP{ // ------------------------ CUBLAS BLAS - like extension (cublas_api.h) // GEAM - {"cublasSgeam", {"hipblasSgeam", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasDgeam", {"hipblasDgeam", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasSgeam", {"hipblasSgeam", CONV_MATH_FUNC, API_BLAS}}, + {"cublasDgeam", {"hipblasDgeam", CONV_MATH_FUNC, API_BLAS}}, {"cublasCgeam", {"hipblasCgeam", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZgeam", {"hipblasZgeam", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2464,8 +2484,8 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasGetVersion_v2", {"hipblasGetVersion", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasSetStream_v2", {"hipblasSetStream", CONV_MATH_FUNC, API_BLAS}}, {"cublasGetStream_v2", {"hipblasGetStream", CONV_MATH_FUNC, API_BLAS}}, - {"cublasGetPointerMode_v2", {"hipblasGetPointerMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasSetPointerMode_v2", {"hipblasSetPointerMode", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasGetPointerMode_v2", {"hipblasGetPointerMode", CONV_MATH_FUNC, API_BLAS}}, + {"cublasSetPointerMode_v2", {"hipblasSetPointerMode", CONV_MATH_FUNC, API_BLAS}}, // GEMV {"cublasSgemv_v2", {"hipblasSgemv", CONV_MATH_FUNC, API_BLAS}}, @@ -2537,7 +2557,7 @@ const std::map CUDA_IDENTIFIER_MAP{ // GER {"cublasSger_v2", {"hipblasSger", CONV_MATH_FUNC, API_BLAS}}, - {"cublasDger_v2", {"hipblasDger", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasDger_v2", {"hipblasDger", CONV_MATH_FUNC, API_BLAS}}, {"cublasCgeru_v2", {"hipblasCgeru", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasCgerc_v2", {"hipblasCgerc", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZgeru_v2", {"hipblasZgeru", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2574,13 +2594,22 @@ const std::map CUDA_IDENTIFIER_MAP{ // Blas3 (v2) Routines // GEMM {"cublasSgemm_v2", {"hipblasSgemm", CONV_MATH_FUNC, API_BLAS}}, - {"cublasDgemm_v2", {"hipblasDgemm", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasDgemm_v2", {"hipblasDgemm", CONV_MATH_FUNC, API_BLAS}}, {"cublasCgemm_v2", {"hipblasCgemm", CONV_MATH_FUNC, API_BLAS}}, + {"cublasCgemm3m", {"hipblasCgemm3m", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasCgemm3mEx", {"hipblasCgemm3mEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasZgemm_v2", {"hipblasZgemm", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasZgemm3m", {"hipblasZgemm3m", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, //IO in FP16 / FP32, computation in float {"cublasSgemmEx", {"hipblasSgemmEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasGemmEx", {"hipblasGemmEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // IO in Int8 complex/cuComplex, computation in cuComplex + {"cublasCgemmEx", {"hipblasCgemmEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + + {"cublasUint8gemmBias", {"hipblasUint8gemmBias", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // SYRK {"cublasSsyrk_v2", {"hipblasSsyrk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2588,8 +2617,17 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasCsyrk_v2", {"hipblasCsyrk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZsyrk_v2", {"hipblasZsyrk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // IO in Int8 complex/cuComplex, computation in cuComplex + {"cublasCsyrkEx", {"hipblasCsyrkEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // IO in Int8 complex/cuComplex, computation in cuComplex, Gaussian math + {"cublasCsyrk3mEx", {"hipblasCsyrk3mEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // HERK {"cublasCherk_v2", {"hipblasCherk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // IO in Int8 complex/cuComplex, computation in cuComplex + {"cublasCherkEx", {"hipblasCherkEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + // IO in Int8 complex/cuComplex, computation in cuComplex, Gaussian math + {"cublasCherk3mEx", {"hipblasCherk3mEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasZherk_v2", {"hipblasZherk", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // SYR2K @@ -2631,6 +2669,9 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasDznrm2_v2", {"hipblasDznrm2", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // DOT + {"cublasDotEx", {"hipblasDotEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasDotcEx", {"hipblasDotcEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasSdot_v2", {"hipblasSdot", CONV_MATH_FUNC, API_BLAS}}, {"cublasDdot_v2", {"hipblasDdot", CONV_MATH_FUNC, API_BLAS}}, @@ -2640,6 +2681,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasZdotc_v2", {"hipblasZdotc", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // SCAL + {"cublasScalEx", {"hipblasScalEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasSscal_v2", {"hipblasSscal", CONV_MATH_FUNC, API_BLAS}}, {"cublasDscal_v2", {"hipblasDscal", CONV_MATH_FUNC, API_BLAS}}, {"cublasCscal_v2", {"hipblasCscal", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2648,6 +2690,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasZdscal_v2", {"hipblasZdscal", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // AXPY + {"cublasAxpyEx", {"hipblasAxpyEx", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasSaxpy_v2", {"hipblasSaxpy", CONV_MATH_FUNC, API_BLAS}}, {"cublasDaxpy_v2", {"hipblasDaxpy", CONV_MATH_FUNC, API_BLAS}}, {"cublasCaxpy_v2", {"hipblasCaxpy", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, @@ -2666,8 +2709,8 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cublasZswap_v2", {"hipblasZswap", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, // AMAX - {"cublasIsamax_v2", {"hipblasIsamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasIdamax_v2", {"hipblasIdamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + {"cublasIsamax_v2", {"hipblasIsamax", CONV_MATH_FUNC, API_BLAS}}, + {"cublasIdamax_v2", {"hipblasIdamax", CONV_MATH_FUNC, API_BLAS}}, {"cublasIcamax_v2", {"hipblasIcamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, {"cublasIzamax_v2", {"hipblasIzamax", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, From 974f8022b1973f7e515b57a00a707003f44ac472 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 21:43:18 +0300 Subject: [PATCH 29/37] [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 30/37] [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 31/37] [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 32/37] [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 33/37] [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 34/37] 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 35/37] 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 36/37] 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 37/37] 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