From 9f62283bad4b79d3af2b8245cbbd2fff8d3ca83a 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 --- tests/src/p2p/hipPeerToPeer_simple.cpp | 42 ++++++++++++-------------- 1 file changed, 20 insertions(+), 22 deletions(-) diff --git a/tests/src/p2p/hipPeerToPeer_simple.cpp b/tests/src/p2p/hipPeerToPeer_simple.cpp index 32cab371b8..c279658b5a 100644 --- a/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/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 bd60ef9031814d3a5e60be0c318e3f97fc4343dd 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 --- tests/src/deviceLib/hip_anyall.cpp | 2 +- tests/src/deviceLib/hip_ballot.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/src/deviceLib/hip_anyall.cpp b/tests/src/deviceLib/hip_anyall.cpp index 06354383df..9d455d15b4 100644 --- a/tests/src/deviceLib/hip_anyall.cpp +++ b/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/tests/src/deviceLib/hip_ballot.cpp b/tests/src/deviceLib/hip_ballot.cpp index 14b8f314a1..e4d3cc70a6 100644 --- a/tests/src/deviceLib/hip_ballot.cpp +++ b/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 3d5b2ef5f68f6824c4b98050aa00dcfeb9d3ffc8 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 --- tests/hipify-clang/lit.cfg | 17 +++++++++++++---- tests/hipify-clang/lit.site.cfg.in | 13 +++++++++++++ tests/hipify-clang/run_test.bat | 17 +++++++++++++++++ 3 files changed, 43 insertions(+), 4 deletions(-) create mode 100644 tests/hipify-clang/run_test.bat diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 76b2ca08dc..4dca50c052 100644 --- a/tests/hipify-clang/lit.cfg +++ b/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/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in index c1f6804d4d..c1095f65a9 100644 --- a/tests/hipify-clang/lit.site.cfg.in +++ b/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/tests/hipify-clang/run_test.bat b/tests/hipify-clang/run_test.bat new file mode 100644 index 0000000000..5db0cc5043 --- /dev/null +++ b/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 e9f94561660dfb6f43f4da28167a91445a9eca5a 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 --- tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 2919939694..0b965a3ae3 100644 --- a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/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 fdc10f085650425fe2a8a25bf22f218489b8f62b 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 --- tests/hipify-clang/run_test.bat | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/hipify-clang/run_test.bat b/tests/hipify-clang/run_test.bat index 5db0cc5043..6eefb7e46e 100644 --- a/tests/hipify-clang/run_test.bat +++ b/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 ebec2903518151e279dfc8e5ff6d1a624a76b889 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 --- hipify-clang/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 79d181528b..2b725d7188 100644 --- a/hipify-clang/CMakeLists.txt +++ b/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 5df3a9bb3b3c8f6cced6658ff9be495ca2d70c58 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. --- tests/hipify-clang/lit.site.cfg.in | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in index c1095f65a9..e52db7b026 100644 --- a/tests/hipify-clang/lit.site.cfg.in +++ b/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 caf6b5c4f07e702e9713c3209851eb480ca51306 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 --- hipify-clang/CMakeLists.txt | 23 +++++++++++++++++++++-- 1 file changed, 21 insertions(+), 2 deletions(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 79d181528b..5910698e51 100644 --- a/hipify-clang/CMakeLists.txt +++ b/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 d9572d7fe7bb3ee556162138d46af09a6024e60e 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 --- docs/markdown/hip_debugging.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md index e7e058d17a..bf877d894e 100644 --- a/docs/markdown/hip_debugging.md +++ b/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 6500cce6aaaf691f6bb9c5e1b805b112b687ceb8 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 --- hipify-clang/CMakeLists.txt | 3 --- tests/hipify-clang/run_test.sh | 11 ----------- 2 files changed, 14 deletions(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 5910698e51..8b3fa7e591 100644 --- a/hipify-clang/CMakeLists.txt +++ b/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/tests/hipify-clang/run_test.sh b/tests/hipify-clang/run_test.sh index 46b2fc066b..418df5dd4d 100755 --- a/tests/hipify-clang/run_test.sh +++ b/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 09655a08530824812349103c377b2d4fba995833 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. --- hipify-clang/src/HipifyAction.cpp | 51 ++++++++++++++++++++------- hipify-clang/src/HipifyAction.h | 9 +++++ tests/hipify-clang/headers_test_03.cu | 10 ++++++ tests/hipify-clang/headers_test_04.cu | 12 +++++++ tests/hipify-clang/headers_test_05.cu | 12 +++++++ 5 files changed, 81 insertions(+), 13 deletions(-) create mode 100644 tests/hipify-clang/headers_test_03.cu create mode 100644 tests/hipify-clang/headers_test_04.cu create mode 100644 tests/hipify-clang/headers_test_05.cu diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 8fb318776d..ada0adec5d 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/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/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index 03d34601f3..a269a37117 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/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/tests/hipify-clang/headers_test_03.cu b/tests/hipify-clang/headers_test_03.cu new file mode 100644 index 0000000000..0223f11119 --- /dev/null +++ b/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/tests/hipify-clang/headers_test_04.cu b/tests/hipify-clang/headers_test_04.cu new file mode 100644 index 0000000000..f10c0eda56 --- /dev/null +++ b/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/tests/hipify-clang/headers_test_05.cu b/tests/hipify-clang/headers_test_05.cu new file mode 100644 index 0000000000..c9428b62d5 --- /dev/null +++ b/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 55fc302e3223c6a3c995143b14a2b7faa89307c5 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 --- Jenkinsfile | 6 +++--- docker/dockerfile-build-ubuntu-16.04 | 3 +++ 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 12c9755103..ce6fbc2e9e 100644 --- a/Jenkinsfile +++ b/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/docker/dockerfile-build-ubuntu-16.04 b/docker/dockerfile-build-ubuntu-16.04 index 031bf72437..8f655f7c78 100644 --- a/docker/dockerfile-build-ubuntu-16.04 +++ b/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 a9d34afc50f189af8fab23e0e56b099449397e35 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 --- Jenkinsfile | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index ce6fbc2e9e..4909e666f0 100644 --- a/Jenkinsfile +++ b/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 e54d9f3df04971b5159a89a830726956500c8369 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 --- hipify-clang/src/HipifyAction.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index ada0adec5d..7cd5b3d402 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/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 58faf1ddc9d15f1e447e88eb775661f438ed58f7 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 --- tests/hipify-clang/lit.cfg | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 4dca50c052..12b1410cee 100644 --- a/tests/hipify-clang/lit.cfg +++ b/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 74b489b5b2fcaad1f9f43bd8c3e02ec5b045f938 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. --- tests/hipify-clang/axpy.cu | 1 - tests/hipify-clang/cudaRegister.cu | 2 -- tests/hipify-clang/square.cu | 2 -- 3 files changed, 5 deletions(-) diff --git a/tests/hipify-clang/axpy.cu b/tests/hipify-clang/axpy.cu index 071c503a35..2e59fc021a 100644 --- a/tests/hipify-clang/axpy.cu +++ b/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/tests/hipify-clang/cudaRegister.cu b/tests/hipify-clang/cudaRegister.cu index 80d17f65b9..79d21707c2 100644 --- a/tests/hipify-clang/cudaRegister.cu +++ b/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 eb61038736bea9f9cc992bd63f2119b476a88d65 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 --- hipify-clang/src/CUDA2HipMap.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 1893f6ce5b..db95c7216b 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 a784a9c824563bd1324b8f877e2fab6c56038c06 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 --- tests/hipify-clang/headers_test_04.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/hipify-clang/headers_test_04.cu b/tests/hipify-clang/headers_test_04.cu index f10c0eda56..57667b5a34 100644 --- a/tests/hipify-clang/headers_test_04.cu +++ b/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 c328e3987cd1484821effe62b89a34357a94a4ac 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 --- tests/hipify-clang/headers_test_03.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/hipify-clang/headers_test_03.cu b/tests/hipify-clang/headers_test_03.cu index 0223f11119..5f2e479683 100644 --- a/tests/hipify-clang/headers_test_03.cu +++ b/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 9c0dc5cfe78b2fba5de5a85a96c391c11203e5c4 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. --- tests/hipify-clang/concurentKernels.cu | 242 ------------------------- 1 file changed, 242 deletions(-) delete mode 100644 tests/hipify-clang/concurentKernels.cu diff --git a/tests/hipify-clang/concurentKernels.cu b/tests/hipify-clang/concurentKernels.cu deleted file mode 100644 index 27e9e0d0e7..0000000000 --- a/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 478fed74fec0232cdac46ed8cc8aad2dc0baf9e8 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 --- ...A_Driver_API_functions_supported_by_HIP.md | 24 +++++++++---------- hipify-clang/src/CUDA2HipMap.cpp | 24 +++++++++---------- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 9a4700b19c..5045c9d81d 100644 --- a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/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/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index db95c7216b..4f468e7b47 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 c0bb1d818fa5c801a7dc995083b183c3017bb44d 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 --- bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/bin/hipify-perl +++ b/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 356eaa81b872d4ca3591ebcdcfaefb0647c28d23 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 --- bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/bin/hipify-perl +++ b/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 5788ac5d379feb1ddcd88bcd5d91298235931342 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) --- ...A_Driver_API_functions_supported_by_HIP.md | 88 +++++++++--------- hipify-clang/src/CUDA2HipMap.cpp | 90 +++++++++---------- 2 files changed, 89 insertions(+), 89 deletions(-) diff --git a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 5045c9d81d..c1c30ae019 100644 --- a/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/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/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..24c97258cf 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 ff5f964c079030bc1398a274d9751ad242562db4 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 --- docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md | 2 +- hipify-clang/src/CUDA2HipMap.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index f27f1538b7..84fd0f469d 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/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/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..b0c4d83706 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 8ff99eeadce5b7062531c6eaffbbabe3fab9cd95 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) --- docs/markdown/CURAND_API_supported_by_HIP.md | 132 +++++++++++++++++++ hipify-clang/src/CUDA2HipMap.cpp | 124 ++++++++++++++++- hipify-clang/src/Statistics.cpp | 4 +- hipify-clang/src/Statistics.h | 2 + 4 files changed, 259 insertions(+), 3 deletions(-) create mode 100644 docs/markdown/CURAND_API_supported_by_HIP.md diff --git a/docs/markdown/CURAND_API_supported_by_HIP.md b/docs/markdown/CURAND_API_supported_by_HIP.md new file mode 100644 index 0000000000..83c3d4a234 --- /dev/null +++ b/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/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 755ce365af..b410ec43f9 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index 00b8a66f11..e2b3632066 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/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/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index da4e296db0..f160ca7383 100644 --- a/hipify-clang/src/Statistics.h +++ b/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 6e000adde4b6bbd06e17cf9da54ec640d417890a 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) --- docs/markdown/CURAND_API_supported_by_HIP.md | 41 ++++++++++++++++++-- hipify-clang/src/CUDA2HipMap.cpp | 37 +++++++++++++++++- 2 files changed, 74 insertions(+), 4 deletions(-) diff --git a/docs/markdown/CURAND_API_supported_by_HIP.md b/docs/markdown/CURAND_API_supported_by_HIP.md index 83c3d4a234..900cfadc03 100644 --- a/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/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/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index b410ec43f9..db78660121 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 ddfb1100808591933e66fe24d41098e00a64fe9d 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 --- hipify-clang/src/CUDA2HipMap.cpp | 97 +++++++++++++++++++++++--------- 1 file changed, 70 insertions(+), 27 deletions(-) diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index db78660121..1f9c6287ed 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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 368db8bf6c66333b989d961b26930fb3da542376 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. --- tests/hipify-clang/run_test.bat | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/tests/hipify-clang/run_test.bat b/tests/hipify-clang/run_test.bat index 6eefb7e46e..b4858db807 100644 --- a/tests/hipify-clang/run_test.bat +++ b/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 c528f4f9c5101f4350c7f31dc141f32d3d8cacaf 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 --- tests/hipify-clang/run_test.bat | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/hipify-clang/run_test.bat b/tests/hipify-clang/run_test.bat index b4858db807..d8c8d74cf0 100644 --- a/tests/hipify-clang/run_test.bat +++ b/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 77f807b597d788d3cc05752175790261d5697c49 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. --- hipify-clang/src/HipifyAction.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..7d9ddebf51 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/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 600d5d7c06f622d6f895fe79fe950f8d0453b4a8 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 --- hipify-clang/src/CUDA2HipMap.cpp | 8 +++--- hipify-clang/src/HipifyAction.cpp | 35 +++++++++++++++++++-------- hipify-clang/src/HipifyAction.h | 5 ++-- tests/hipify-clang/headers_test_01.cu | 4 ++- tests/hipify-clang/headers_test_02.cu | 10 +++++--- tests/hipify-clang/headers_test_06.cu | 8 ++++++ tests/hipify-clang/headers_test_07.cu | 8 ++++++ tests/hipify-clang/headers_test_08.cu | 14 +++++++++++ 8 files changed, 71 insertions(+), 21 deletions(-) create mode 100644 tests/hipify-clang/headers_test_06.cu create mode 100644 tests/hipify-clang/headers_test_07.cu create mode 100644 tests/hipify-clang/headers_test_08.cu diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 1f9c6287ed..6bcc0b38c4 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/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/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..9ac0eacbd1 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/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/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index a269a37117..42622c1e01 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/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/tests/hipify-clang/headers_test_01.cu b/tests/hipify-clang/headers_test_01.cu index c39ef80d8f..3747c339e8 100644 --- a/tests/hipify-clang/headers_test_01.cu +++ b/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/tests/hipify-clang/headers_test_02.cu b/tests/hipify-clang/headers_test_02.cu index 90d412f797..57308efd59 100644 --- a/tests/hipify-clang/headers_test_02.cu +++ b/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/tests/hipify-clang/headers_test_06.cu b/tests/hipify-clang/headers_test_06.cu new file mode 100644 index 0000000000..bce73c42df --- /dev/null +++ b/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/tests/hipify-clang/headers_test_07.cu b/tests/hipify-clang/headers_test_07.cu new file mode 100644 index 0000000000..4237e1eb72 --- /dev/null +++ b/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/tests/hipify-clang/headers_test_08.cu b/tests/hipify-clang/headers_test_08.cu new file mode 100644 index 0000000000..ad54871bd8 --- /dev/null +++ b/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 f9416a0c49737f0ac19f42565ab6e93a5ea178f1 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 --- tests/hipify-clang/cudaRegister.cu | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/tests/hipify-clang/cudaRegister.cu b/tests/hipify-clang/cudaRegister.cu index 79d21707c2..43b4345337 100644 --- a/tests/hipify-clang/cudaRegister.cu +++ b/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 0497424978de459d817432efbafb73c2df544b46 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 --- Jenkinsfile | 3 ++- include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 4909e666f0..6142b94d95 100644 --- a/Jenkinsfile +++ b/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/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..2bec0017db 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/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 8acc8365fa9a424e76c5b1d6d34b121c23ed8de9 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 --- include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..1b2dfc1c7c 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/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 98f3fe3939148599367f24faf27564ffdf926230 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 --- Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 6142b94d95..6e5f7bc8e5 100644 --- a/Jenkinsfile +++ b/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 9860d5de20116d88943ed6512ebd9ea95d87a22b 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 --- packaging/hip_doc.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/packaging/hip_doc.txt b/packaging/hip_doc.txt index daef7810b4..5987c2c128 100644 --- a/packaging/hip_doc.txt +++ b/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