From 3c6e6c100ffcffa218d8fb7280bf2a5f80f17f70 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 20 Dec 2017 09:36:00 +0530 Subject: [PATCH 01/37] Return pass on single gpu in hipPeerToPeer_simple [ROCm/clr commit: 037ce74fc94cf55e7c5ed8fb2d81b7fc8988ba5b] --- .../tests/src/p2p/hipPeerToPeer_simple.cpp | 42 +++++++++---------- 1 file changed, 20 insertions(+), 22 deletions(-) diff --git a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp index 32cab371b8..c279658b5a 100644 --- a/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/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 bab8090dfc14b7c49e284b928e04a458e56886cf Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 20 Dec 2017 12:05:21 +0530 Subject: [PATCH 02/37] [nvccWarnings] Fix -Wno-deprecated-declarations in hip_anyall and hip_ballot [ROCm/clr commit: f69762b300aef58f3d60f45ef2575793a824efcc] --- projects/clr/hipamd/tests/src/deviceLib/hip_anyall.cpp | 2 +- projects/clr/hipamd/tests/src/deviceLib/hip_ballot.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hip_anyall.cpp b/projects/clr/hipamd/tests/src/deviceLib/hip_anyall.cpp index 06354383df..9d455d15b4 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hip_anyall.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hip_anyall.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations * RUN: %t * HIT_END */ diff --git a/projects/clr/hipamd/tests/src/deviceLib/hip_ballot.cpp b/projects/clr/hipamd/tests/src/deviceLib/hip_ballot.cpp index 14b8f314a1..e4d3cc70a6 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hip_ballot.cpp +++ b/projects/clr/hipamd/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 3e87dc78e6bbdf25ba639766a3fb44c7f42e9cb2 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 9 Jan 2018 20:20:28 +0300 Subject: [PATCH 03/37] [HIPIFY][tests] Add Windows testing support [ROCm/clr commit: 257bc4748cb7ef485be1a7bfe8e8967baac74a5d] --- projects/clr/hipamd/tests/hipify-clang/lit.cfg | 17 +++++++++++++---- .../hipamd/tests/hipify-clang/lit.site.cfg.in | 13 +++++++++++++ .../clr/hipamd/tests/hipify-clang/run_test.bat | 17 +++++++++++++++++ 3 files changed, 43 insertions(+), 4 deletions(-) create mode 100644 projects/clr/hipamd/tests/hipify-clang/run_test.bat diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index 76b2ca08dc..4dca50c052 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/tests/hipify-clang/lit.cfg @@ -44,8 +44,17 @@ if obj_root is not None: path = os.path.pathsep.join((llvm_tools_dir, config.environment['PATH'])) config.environment['PATH'] = path -config.substitutions.append(("hipify", obj_root+"/hipify-clang")) +hipify_path = obj_root +clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'" -# Clang args for CUDA... -config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30 -isystem%s/samples/common/inc" % (config.cuda_root, config.cuda_root))) -config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) +if sys.platform in ['win32']: + run_test_ext = ".bat" + hipify_path += "/" + config.build_type + clang_args += " -isystem'%s'/common/inc -std=c++14" +else: + run_test_ext = ".sh" + clang_args += " -isystem'%s'/samples/common/inc" + +config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root))) +config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"')) +config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"')) diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in index c1f6804d4d..c1095f65a9 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in +++ b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in @@ -1,8 +1,21 @@ import sys +import os config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" +if sys.platform in ['win32']: + config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@" + if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND": + config.cuda_samples_root = os.environ.get('NVCUDASAMPLES_ROOT') + if not config.cuda_samples_root or config.cuda_samples_root == "NVCUDASAMPLES_ROOT-NOTFOUND": + lit_config.fatal('No CUDA Samples dir set! Please set CUDA_SDK_ROOT_DIR.') + config.cuda_sdk_root = config.cuda_samples_root + config.build_type = "@CMAKE_BUILD_TYPE@" + if not config.build_type: + config.build_type = "Debug" +else: + config.cuda_sdk_root = config.cuda_root # Support substitution of the tools and libs dirs with user parameters. This is # used when we can't determine the tool dir at configuration time. diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.bat b/projects/clr/hipamd/tests/hipify-clang/run_test.bat new file mode 100644 index 0000000000..5db0cc5043 --- /dev/null +++ b/projects/clr/hipamd/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 015366e29ed91deb18e1ca92978dbc82ad2ad779 Mon Sep 17 00:00:00 2001 From: Phaneendr-kumar Lanka Date: Wed, 10 Jan 2018 10:51:01 +0530 Subject: [PATCH 04/37] [nvcc] Enable hipGetDeviceAttribute [ROCm/clr commit: dc6094cc608f846325c84efc6f0b9c89839c85d5] --- .../tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 2919939694..0b965a3ae3 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/clr/hipamd/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 9a0d16f791e093a9744828838f307aaf997a1737 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 10 Jan 2018 21:03:02 +0300 Subject: [PATCH 05/37] [HIPIFY][tests] Add setlocal to batch script [ROCm/clr commit: b32639d1a8bf8050a3555d0c5daeae2cc210b246] --- projects/clr/hipamd/tests/hipify-clang/run_test.bat | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.bat b/projects/clr/hipamd/tests/hipify-clang/run_test.bat index 5db0cc5043..6eefb7e46e 100644 --- a/projects/clr/hipamd/tests/hipify-clang/run_test.bat +++ b/projects/clr/hipamd/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 b067e12c7c91685792dae816847bd50fc0a4ed8f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 10 Jan 2018 21:06:06 +0300 Subject: [PATCH 06/37] [HIPIFY][cmake] Exclude socat from Win config [ROCm/clr commit: 5c15cc77cb087e5907d5704816522c1043c8d358] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 79d181528b..2b725d7188 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/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 674bcc42a3f1588ec4c4d22c0f96f71de8b8984e Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 12 Jan 2018 17:15:37 +0300 Subject: [PATCH 07/37] [HIPIFY][tests][win] CUDA samples root env. var is changes Env. var NVCUDASAMPLES_ROOT is changed to NVCUDASAMPLESX_Y_ROOT where X - major ver, Y - minor ver. Reason: NVCUDASAMPLES_ROOT contains path to CUDA SDK installed last, while NVCUDASAMPLESX_Y_ROOT contains samples of the same version as of CUDA_TOOLKIT_ROOT_DIR. [ROCm/clr commit: 39a03720770205bd4fedd7f5d674f4ac08dae734] --- projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in index c1095f65a9..e52db7b026 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in +++ b/projects/clr/hipamd/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 ec693ff0d741d8168e994f5d033c448d4ac6e4c7 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 12 Jan 2018 17:36:41 +0300 Subject: [PATCH 08/37] [HIPIFY][cmake] Version compatibility checks of CUDA and clang are added [ROCm/clr commit: d11dccdd443c3295f7f02fdb659c68b7558ae717] --- .../clr/hipamd/hipify-clang/CMakeLists.txt | 23 +++++++++++++++++-- 1 file changed, 21 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 79d181528b..5910698e51 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/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 a880e5918d2352d4a52fbcb8d655a799bddee32a Mon Sep 17 00:00:00 2001 From: Peng Date: Sat, 13 Jan 2018 12:33:27 -0600 Subject: [PATCH 09/37] Update hip_debugging.md on typo of chicken bits [ROCm/clr commit: 0831736e70f59af73ede6b5a943d03d208df15ec] --- projects/clr/hipamd/docs/markdown/hip_debugging.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_debugging.md b/projects/clr/hipamd/docs/markdown/hip_debugging.md index e7e058d17a..bf877d894e 100644 --- a/projects/clr/hipamd/docs/markdown/hip_debugging.md +++ b/projects/clr/hipamd/docs/markdown/hip_debugging.md @@ -74,7 +74,7 @@ HIP provides 3 environment variables in the HIP_*_BLOCKING family. These introd - HIP_API_BLOCKING : Forces hipMemcpyAsync and hipMemsetAsync to be host-synchronous, meaning they will wait for the requested operation to complete before returning to the caller. These options cause HCC to serialize. Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP. -- HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize. +- HCC_SERIALIZE_KERNEL : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize. - HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize. - HSA_ENABLE_SDMA=0 : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine. This flag is useful to isolate issues with the hardware copy engines. From ed9caac5875f567a672c28fa326870307cc960e3 Mon Sep 17 00:00:00 2001 From: emankov Date: Mon, 15 Jan 2018 14:20:37 +0300 Subject: [PATCH 10/37] [HIPIFY][#311][fix] Get rid of socat in run_test.sh [ROCm/clr commit: f83df46b8cf99008a2637b3f45db9f7b4a876818] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 3 --- projects/clr/hipamd/tests/hipify-clang/run_test.sh | 11 ----------- 2 files changed, 14 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 5910698e51..8b3fa7e591 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -78,9 +78,6 @@ if (HIPIFY_CLANG_TESTS) require_program(lit) require_program(FileCheck) - if(NOT WIN32) - require_program(socat) - endif() # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the # value of --cuda-path for the test runs. diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.sh b/projects/clr/hipamd/tests/hipify-clang/run_test.sh index 46b2fc066b..418df5dd4d 100755 --- a/projects/clr/hipamd/tests/hipify-clang/run_test.sh +++ b/projects/clr/hipamd/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 36b0e4003ebbf1266bcc2dda14b5cc69824a3022 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 15 Jan 2018 21:05:05 +0300 Subject: [PATCH 11/37] [HIPIFY][fix][#211] Algorithm for explicit insert of hip include directive If in source CUDA file main header (cuda_runtime.h or cuda.h) is not presented, corresponding HIP main header (hip_runtime.h) should be explicitly included in output hipified file. [Algorithm] 1. If #pragma once is presented, HIP main header should be placed just after it; 2. Otherwise if any other (not CUDA main) header is presented, HIP main header should be placed just before it; 3. Otherwise HIP main header should be placed in the beginning of output file. P.S. There might be one more situation when #ifndef #define ... #endif guard for the entire file is presented (make sense for *.h, *.hpp, *.cuh files). In this case HIP main include should be placed just after such #ifdef, or after #pragma once, if it is also presented. This situation will be handled in a separate change. [ROCm/clr commit: e90a76a1ef6ecd719f3f89778ec60cae3d0a2d2b] --- .../hipamd/hipify-clang/src/HipifyAction.cpp | 51 ++++++++++++++----- .../hipamd/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 projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/headers_test_05.cu diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 8fb318776d..ada0adec5d 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -152,7 +152,10 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const auto found = CUDA_INCLUDE_MAP.find(file_name); if (found == CUDA_INCLUDE_MAP.end()) { - // Not a CUDA include - don't touch it. + if (!firstNotMainHeader) { + firstNotMainHeader = true; + firstNotMainHeaderLoc = hash_loc; + } return; } @@ -160,7 +163,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, bool secondMainInclude = false; if (found->second.hipName == "hip/hip_runtime.h") { if (insertedRuntimeHeader) { - secondMainInclude = true; + secondMainInclude = true; } insertedRuntimeHeader = true; } @@ -178,15 +181,15 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, // Keep the same include type that the user gave. if (!secondMainInclude) { - clang::SmallString<128> includeBuffer; - if (is_angled) { - newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); - } else { - newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); - } + clang::SmallString<128> includeBuffer; + if (is_angled) { + newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); + } else { + newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); + } } else { - // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. - sl = hash_loc; + // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. + sl = hash_loc; } const char *B = SM.getCharacterData(sl); const char *E = SM.getCharacterData(filename_range.getEnd()); @@ -194,6 +197,18 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, insertReplacement(Rep, clang::FullSourceLoc{sl, SM}); } +void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) { + if (pragmaOnce) { return; } + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + const clang::Token tok = PP.LookAhead(0); + clang::LangOptions DefaultLangOptions; + StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); + if (Text == "once") { + pragmaOnce = true; + pragmaOnceLoc = PP.LookAhead(1).getLocation(); + } +} bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) { StringRef refName = "cudaLaunchKernel"; @@ -339,10 +354,16 @@ void HipifyAction::EndSourceFileAction() { // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert // one copy of the hip include into every file. clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - - clang::SourceLocation sl = SM.getLocForStartOfFile(SM.getMainFileID()); + clang::SourceLocation sl; + if (pragmaOnce) { + sl = pragmaOnceLoc; + } else if (firstNotMainHeader) { + sl = firstNotMainHeaderLoc; + } else { + sl = SM.getLocForStartOfFile(SM.getMainFileID()); + } clang::FullSourceLoc fullSL(sl, SM); - ct::Replacement Rep(SM, sl, 0, "#include \n"); + ct::Replacement Rep(SM, sl, 0, "\n#include \n"); insertReplacement(Rep, fullSL); } @@ -367,6 +388,10 @@ public: const clang::Module* imported) override { hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported); } + + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override { + hipifyAction.PragmaDirective(Loc, Introducer); + } }; } diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h index 03d34601f3..a269a37117 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h @@ -23,6 +23,10 @@ private: // not, we insert it at the top of the file when we finish processing it. // This approach means we do the best it's possible to do w.r.t preserving the user's include order. bool insertedRuntimeHeader = false; + bool firstNotMainHeader = false; + bool pragmaOnce = false; + clang::SourceLocation firstNotMainHeaderLoc; + clang::SourceLocation pragmaOnceLoc; /** * Rewrite a string literal to refer to hip, not CUDA. @@ -57,6 +61,11 @@ public: StringRef relative_path, const clang::Module *imported); + /** + * Called by the preprocessor for each pragma directive during the non-raw lexing pass. + */ + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer); + protected: /** * Add a Replacement for the current file. These will all be applied after executing the FrontendAction. diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu new file mode 100644 index 0000000000..0223f11119 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu @@ -0,0 +1,10 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #pragma once +// CHECK-NEXT: #include +#pragma once +// CHECK-NOT: #include +int main(int argc, char* argv[]) { + return 0; +} + diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu new file mode 100644 index 0000000000..f10c0eda56 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu @@ -0,0 +1,12 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NEXT: #include +// CHECK-NEXT: #include +#include +#include +// CHECK-NOT: #include +int main(int argc, char* argv[]) { + return 0; +} + diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_05.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_05.cu new file mode 100644 index 0000000000..c9428b62d5 --- /dev/null +++ b/projects/clr/hipamd/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 99268efb2c3ab556dbd5b8a9074de3ad63cebe2e Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 10 Jan 2018 11:26:51 -0600 Subject: [PATCH 12/37] Adding dependencies for rocm_agent_enumerator [ROCm/clr commit: f51ac2b43eda4355e302c357f81adaf9eacc9abc] --- projects/clr/hipamd/Jenkinsfile | 6 +++--- projects/clr/hipamd/docker/dockerfile-build-ubuntu-16.04 | 3 +++ 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/Jenkinsfile b/projects/clr/hipamd/Jenkinsfile index 12c9755103..ce6fbc2e9e 100644 --- a/projects/clr/hipamd/Jenkinsfile +++ b/projects/clr/hipamd/Jenkinsfile @@ -371,7 +371,7 @@ parallel hcc_ctu: { String hcc_ver = 'hcc-ctu' String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/clang_tot_upgrade/hcc-lc-ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd' + String inside_args = '--device=/dev/kfd --device=/dev/dri' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) @@ -411,8 +411,8 @@ hcc_1_6: node('docker && rocm') { String hcc_ver = 'hcc-1.6' - String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/roc-1.6.x/hcc-lc-ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd' + String from_image = 'rocm/dev-ubuntu-16.04:latest' + String inside_args = '--device=/dev/kfd --device=/dev/dri' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) diff --git a/projects/clr/hipamd/docker/dockerfile-build-ubuntu-16.04 b/projects/clr/hipamd/docker/dockerfile-build-ubuntu-16.04 index 031bf72437..8f655f7c78 100644 --- a/projects/clr/hipamd/docker/dockerfile-build-ubuntu-16.04 +++ b/projects/clr/hipamd/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 adb16194c41b594edc9d2e4bb7a35fef161e468e Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 10 Jan 2018 13:28:18 -0600 Subject: [PATCH 13/37] adding group-add flag to docker run [ROCm/clr commit: e645a2b8f397ef40128ebd151bf1d54d1c13672e] --- projects/clr/hipamd/Jenkinsfile | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/Jenkinsfile b/projects/clr/hipamd/Jenkinsfile index ce6fbc2e9e..4909e666f0 100644 --- a/projects/clr/hipamd/Jenkinsfile +++ b/projects/clr/hipamd/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 38c69ecb73977708f66016ed82bd74ab72f55df9 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 15:08:08 +0300 Subject: [PATCH 14/37] Update HipifyAction.cpp dead code eliminate [ROCm/clr commit: a26e323612c21137698ba820f248e9622128cd04] --- projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index ada0adec5d..7cd5b3d402 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/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 dea3b9ed95b903808b039076eebf4a36d442c889 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 16:40:31 +0300 Subject: [PATCH 15/37] [HIPIFY][tests] Add more suffixes to lit config [ROCm/clr commit: 5c82a2e7fa8d3c2ef8f78c83af575c275db38a49] --- projects/clr/hipamd/tests/hipify-clang/lit.cfg | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index 4dca50c052..12b1410cee 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/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 0769a80c9409a291ca8204914e5df49cdb1ba162 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 17:13:29 +0300 Subject: [PATCH 16/37] [HIPIFY][tests] Remove checks on cudaBuiltins As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim. Other CUDA builtins are not tracked for now. [ROCm/clr commit: 42f0966a9efa7a489d5d4fe501c613b0994f9042] --- projects/clr/hipamd/tests/hipify-clang/axpy.cu | 1 - projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu | 2 -- projects/clr/hipamd/tests/hipify-clang/square.cu | 2 -- 3 files changed, 5 deletions(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index 071c503a35..2e59fc021a 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -16,7 +16,6 @@ template __global__ void axpy(T a, T *x, T *y) { - // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } diff --git a/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu b/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu index 80d17f65b9..79d21707c2 100644 --- a/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu +++ b/projects/clr/hipamd/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 5a369c5d2964ced148143e495fd18458ec7001ec Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 17:21:19 +0300 Subject: [PATCH 17/37] [HIPIFY] Add more supported by HIP CUDA RT API Textures and Arrays data types [ROCm/clr commit: 23889e79872bf382aca5dd18bf9a3e28dbebad08] --- .../clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 1893f6ce5b..db95c7216b 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 f7605413bf7968b65adfddca460558b08acf57f5 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 19:21:14 +0300 Subject: [PATCH 18/37] Update headers_test_04.cu [ROCm/clr commit: 44d51e794bdd3cca5da823174888e7223ab187ab] --- projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu index f10c0eda56..57667b5a34 100644 --- a/projects/clr/hipamd/tests/hipify-clang/headers_test_04.cu +++ b/projects/clr/hipamd/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 adcd58db5c9a1b1c399192a8f0913699db6d8f9d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 19:21:59 +0300 Subject: [PATCH 19/37] Update headers_test_03.cu [ROCm/clr commit: 3db7dc5b9e69ef7e3f5f0cf0d8bb672d521a6b56] --- projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu index 0223f11119..5f2e479683 100644 --- a/projects/clr/hipamd/tests/hipify-clang/headers_test_03.cu +++ b/projects/clr/hipamd/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 9c4549fac4f2fd70e8270a60922efec0750491b3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 20:41:08 +0300 Subject: [PATCH 20/37] [HIPIFY][tests] remove concurentKernels.cu as it is one of CUDA SDK samples. [ROCm/clr commit: 284d1cb4e3580e61682f182433cfe0e094fcf578] --- .../tests/hipify-clang/concurentKernels.cu | 242 ------------------ 1 file changed, 242 deletions(-) delete mode 100644 projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu b/projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu deleted file mode 100644 index 27e9e0d0e7..0000000000 --- a/projects/clr/hipamd/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 c5c57066701dd676923bb38a8353f90cba437697 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 16 Jan 2018 21:07:50 +0300 Subject: [PATCH 21/37] [HIPIFY] Add more supported by HIP CUDA Driver API Arrays data types and functions [ROCm/clr commit: e2ac6c6296fb11462998e48a8f69ab82f0cc1871] --- ...A_Driver_API_functions_supported_by_HIP.md | 24 +++++++++---------- .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 24 +++++++++---------- 2 files changed, 24 insertions(+), 24 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 9a4700b19c..5045c9d81d 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -5,8 +5,8 @@ | **type** | **CUDA** | **HIP** | |-------------:|---------------------------------------------------------------|------------------------------------------------------------| | struct | `CUDA_ARRAY3D_DESCRIPTOR` | | -| struct | `CUDA_ARRAY_DESCRIPTOR` | | -| struct | `CUDA_MEMCPY2D` | | +| struct | `CUDA_ARRAY_DESCRIPTOR` | `HIP_ARRAY_DESCRIPTOR` | +| struct | `CUDA_MEMCPY2D` | `hip_Memcpy2D` | | struct | `CUDA_MEMCPY3D` | | | struct | `CUDA_MEMCPY3D_PEER` | | | struct | `CUDA_POINTER_ATTRIBUTE_P2P_TOKENS` | | @@ -27,15 +27,15 @@ | 0x03 |*`CU_CUBEMAP_FACE_NEGATIVE_Y`* | | | 0x04 |*`CU_CUBEMAP_FACE_POSITIVE_Z`* | | | 0x05 |*`CU_CUBEMAP_FACE_NEGATIVE_Z`* | | -| enum |***`CUarray_format`*** | | -| 0x01 |*`CU_AD_FORMAT_UNSIGNED_INT8`* | | -| 0x02 |*`CU_AD_FORMAT_UNSIGNED_INT16`* | | -| 0x03 |*`CU_AD_FORMAT_UNSIGNED_INT32`* | | -| 0x08 |*`CU_AD_FORMAT_SIGNED_INT8`* | | -| 0x09 |*`CU_AD_FORMAT_SIGNED_INT16`* | | -| 0x0a |*`CU_AD_FORMAT_SIGNED_INT32`* | | -| 0x10 |*`CU_AD_FORMAT_HALF`* | | -| 0x20 |*`CU_AD_FORMAT_FLOAT`* | | +| enum |***`CUarray_format`*** |***`hipArray_format`*** | +| 0x01 |*`CU_AD_FORMAT_UNSIGNED_INT8`* |*`HIP_AD_FORMAT_UNSIGNED_INT8`* | +| 0x02 |*`CU_AD_FORMAT_UNSIGNED_INT16`* |*`HIP_AD_FORMAT_UNSIGNED_INT16`* | +| 0x03 |*`CU_AD_FORMAT_UNSIGNED_INT32`* |*`HIP_AD_FORMAT_UNSIGNED_INT32`* | +| 0x08 |*`CU_AD_FORMAT_SIGNED_INT8`* |*`HIP_AD_FORMAT_SIGNED_INT8`* | +| 0x09 |*`CU_AD_FORMAT_SIGNED_INT16`* |*`HIP_AD_FORMAT_SIGNED_INT16`* | +| 0x0a |*`CU_AD_FORMAT_SIGNED_INT32`* |*`HIP_AD_FORMAT_SIGNED_INT32`* | +| 0x10 |*`CU_AD_FORMAT_HALF`* |*`HIP_AD_FORMAT_HALF`* | +| 0x20 |*`CU_AD_FORMAT_FLOAT`* |*`HIP_AD_FORMAT_FLOAT`* | | enum |***`CUctx_flags`*** | | | 0x00 |*`CU_CTX_SCHED_AUTO`* | | | 0x01 |*`CU_CTX_SCHED_SPIN`* | | @@ -518,7 +518,7 @@ | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------| -| `cuArray3DCreate` | | +| `cuArray3DCreate` | `hipArray3DCreate` | | `cuArray3DGetDescriptor` | | | `cuArrayCreate` | | | `cuArrayDestroy` | | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index db95c7216b..4f468e7b47 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 3872cd6b3fbe1ba94e38f7585b1595052490c60f Mon Sep 17 00:00:00 2001 From: Ashish Date: Wed, 17 Jan 2018 14:39:42 -0800 Subject: [PATCH 22/37] Fixed the mapping of cudaPointerGetAttributes -> hipPointerGetAttributes in the hipify-perl [ROCm/clr commit: e8cd2779716612d37c2da51e466c1319835d63bd] --- projects/clr/hipamd/bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/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 3ed079667fa7b8b931be618ed04b6a05802e73eb Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 09:14:16 +0300 Subject: [PATCH 23/37] [HIPIFY][perl] Typo fix [ROCm/clr commit: 87895b7d01b5a9548a2cec5939e14c7e3aca4d36] --- projects/clr/hipamd/bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index 65e3d26ee0..3ea201d14f 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/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 7bed930cd489ae8a8fdab02a44fc2559c072efb4 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 12:03:03 +0300 Subject: [PATCH 24/37] [HIPIFY] Add CUDA Driver API Texture Ref support (partial) [ROCm/clr commit: 2bf8cc691b5f522717f1e9915cb6e13a8a55eb1f] --- ...A_Driver_API_functions_supported_by_HIP.md | 88 +++++++++--------- .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 90 +++++++++---------- 2 files changed, 89 insertions(+), 89 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 5045c9d81d..c1c30ae019 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -248,42 +248,42 @@ | 0x01 |*`CU_RESOURCE_TYPE_MIPMAPPED_ARRAY`* | | | 0x02 |*`CU_RESOURCE_TYPE_LINEAR`* | | | 0x03 |*`CU_RESOURCE_TYPE_PITCH2D`* | | -| enum |***`CUresourceViewFormat`*** | | -| 0x00 |*`CU_RES_VIEW_FORMAT_NONE`* | | -| 0x01 |*`CU_RES_VIEW_FORMAT_UINT_1X8`* | | -| 0x02 |*`CU_RES_VIEW_FORMAT_UINT_2X8`* | | -| 0x03 |*`CU_RES_VIEW_FORMAT_UINT_4X8`* | | -| 0x04 |*`CU_RES_VIEW_FORMAT_SINT_1X8`* | | -| 0x05 |*`CU_RES_VIEW_FORMAT_SINT_2X8`* | | -| 0x06 |*`CU_RES_VIEW_FORMAT_SINT_4X8`* | | -| 0x07 |*`CU_RES_VIEW_FORMAT_UINT_1X16`* | | -| 0x08 |*`CU_RES_VIEW_FORMAT_UINT_2X16`* | | -| 0x09 |*`CU_RES_VIEW_FORMAT_UINT_4X16`* | | -| 0x0a |*`CU_RES_VIEW_FORMAT_SINT_1X16`* | | -| 0x0b |*`CU_RES_VIEW_FORMAT_SINT_2X16`* | | -| 0x0c |*`CU_RES_VIEW_FORMAT_SINT_4X16`* | | -| 0x0d |*`CU_RES_VIEW_FORMAT_UINT_1X32`* | | -| 0x0e |*`CU_RES_VIEW_FORMAT_UINT_2X32`* | | -| 0x0f |*`CU_RES_VIEW_FORMAT_UINT_4X32`* | | -| 0x10 |*`CU_RES_VIEW_FORMAT_SINT_1X32`* | | -| 0x11 |*`CU_RES_VIEW_FORMAT_SINT_2X32`* | | -| 0x12 |*`CU_RES_VIEW_FORMAT_SINT_4X32`* | | -| 0x13 |*`CU_RES_VIEW_FORMAT_FLOAT_1X16`* | | -| 0x14 |*`CU_RES_VIEW_FORMAT_FLOAT_2X16`* | | -| 0x15 |*`CU_RES_VIEW_FORMAT_FLOAT_4X16`* | | -| 0x16 |*`CU_RES_VIEW_FORMAT_FLOAT_1X32`* | | -| 0x17 |*`CU_RES_VIEW_FORMAT_FLOAT_2X32`* | | -| 0x18 |*`CU_RES_VIEW_FORMAT_FLOAT_4X32`* | | -| 0x19 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC1`* | | -| 0x1a |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* | | -| 0x1b |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* | | -| 0x1c |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC4`* | | -| 0x1d |*`CU_RES_VIEW_FORMAT_SIGNED_BC4`* | | -| 0x1e |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC5`* | | -| 0x1f |*`CU_RES_VIEW_FORMAT_SIGNED_BC5`* | | -| 0x20 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC6H`* | | -| 0x21 |*`CU_RES_VIEW_FORMAT_SIGNED_BC6H`* | | -| 0x22 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC7`* | | +| enum |***`CUresourceViewFormat`*** |***`hipResourceViewFormat`*** | +| 0x00 |*`CU_RES_VIEW_FORMAT_NONE`* |*`hipResViewFormatNone`* | +| 0x01 |*`CU_RES_VIEW_FORMAT_UINT_1X8`* |*`hipResViewFormatUnsignedChar1`* | +| 0x02 |*`CU_RES_VIEW_FORMAT_UINT_2X8`* |*`hipResViewFormatUnsignedChar2`* | +| 0x03 |*`CU_RES_VIEW_FORMAT_UINT_4X8`* |*`hipResViewFormatUnsignedChar4`* | +| 0x04 |*`CU_RES_VIEW_FORMAT_SINT_1X8`* |*`hipResViewFormatSignedChar1`* | +| 0x05 |*`CU_RES_VIEW_FORMAT_SINT_2X8`* |*`hipResViewFormatSignedChar2`* | +| 0x06 |*`CU_RES_VIEW_FORMAT_SINT_4X8`* |*`hipResViewFormatSignedChar4`* | +| 0x07 |*`CU_RES_VIEW_FORMAT_UINT_1X16`* |*`hipResViewFormatUnsignedShort1`* | +| 0x08 |*`CU_RES_VIEW_FORMAT_UINT_2X16`* |*`hipResViewFormatUnsignedShort2`* | +| 0x09 |*`CU_RES_VIEW_FORMAT_UINT_4X16`* |*`hipResViewFormatUnsignedShort4`* | +| 0x0a |*`CU_RES_VIEW_FORMAT_SINT_1X16`* |*`hipResViewFormatSignedShort1`* | +| 0x0b |*`CU_RES_VIEW_FORMAT_SINT_2X16`* |*`hipResViewFormatSignedShort2`* | +| 0x0c |*`CU_RES_VIEW_FORMAT_SINT_4X16`* |*`hipResViewFormatSignedShort4`* | +| 0x0d |*`CU_RES_VIEW_FORMAT_UINT_1X32`* |*`hipResViewFormatUnsignedInt1`* | +| 0x0e |*`CU_RES_VIEW_FORMAT_UINT_2X32`* |*`hipResViewFormatUnsignedInt2`* | +| 0x0f |*`CU_RES_VIEW_FORMAT_UINT_4X32`* |*`hipResViewFormatUnsignedInt4`* | +| 0x10 |*`CU_RES_VIEW_FORMAT_SINT_1X32`* |*`hipResViewFormatSignedInt1`* | +| 0x11 |*`CU_RES_VIEW_FORMAT_SINT_2X32`* |*`hipResViewFormatSignedInt2`* | +| 0x12 |*`CU_RES_VIEW_FORMAT_SINT_4X32`* |*`hipResViewFormatSignedInt4`* | +| 0x13 |*`CU_RES_VIEW_FORMAT_FLOAT_1X16`* |*`hipResViewFormatHalf1`* | +| 0x14 |*`CU_RES_VIEW_FORMAT_FLOAT_2X16`* |*`hipResViewFormatHalf2`* | +| 0x15 |*`CU_RES_VIEW_FORMAT_FLOAT_4X16`* |*`hipResViewFormatHalf4`* | +| 0x16 |*`CU_RES_VIEW_FORMAT_FLOAT_1X32`* |*`hipResViewFormatFloat1`* | +| 0x17 |*`CU_RES_VIEW_FORMAT_FLOAT_2X32`* |*`hipResViewFormatFloat2`* | +| 0x18 |*`CU_RES_VIEW_FORMAT_FLOAT_4X32`* |*`hipResViewFormatFloat4`* | +| 0x19 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC1`* |*`hipResViewFormatUnsignedBlockCompressed1`* | +| 0x1a |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC2`* |*`hipResViewFormatUnsignedBlockCompressed2`* | +| 0x1b |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC3`* |*`hipResViewFormatUnsignedBlockCompressed3`* | +| 0x1c |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC4`* |*`hipResViewFormatUnsignedBlockCompressed4`* | +| 0x1d |*`CU_RES_VIEW_FORMAT_SIGNED_BC4`* |*`hipResViewFormatSignedBlockCompressed4`* | +| 0x1e |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC5`* |*`hipResViewFormatUnsignedBlockCompressed5`* | +| 0x1f |*`CU_RES_VIEW_FORMAT_SIGNED_BC5`* |*`hipResViewFormatSignedBlockCompressed5`* | +| 0x20 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC6H`* |*`hipResViewFormatUnsignedBlockCompressed6H`* | +| 0x21 |*`CU_RES_VIEW_FORMAT_SIGNED_BC6H`* |*`hipResViewFormatSignedBlockCompressed6H`* | +| 0x22 |*`CU_RES_VIEW_FORMAT_UNSIGNED_BC7`* |*`hipResViewFormatUnsignedBlockCompressed7`* | | enum |***`CUresult`*** |***`hipError_t`*** | | 0 |*`CUDA_SUCCESS`* |*`hipSuccess`* | | 1 |*`CUDA_ERROR_INVALID_VALUE`* |*`hipErrorInvalidValue`* | @@ -394,8 +394,8 @@ | typedef | `CUstreamCallback` | `hipStreamCallback_t` | | typedef | `CUsurfObject` | | | typedef | `CUsurfref` | | -| typedef | `CUtexObject` | | -| typedef | `CUtexref` | | +| typedef | `CUtexObject` | `hipTextureObject_t` | +| typedef | `CUtexref` | `textureReference` | | define |`CU_IPC_HANDLE_SIZE` | | | define |`CU_LAUNCH_PARAM_BUFFER_POINTER` | `HIP_LAUNCH_PARAM_BUFFER_POINTER` | | define |`CU_LAUNCH_PARAM_BUFFER_SIZE` | `HIP_LAUNCH_PARAM_BUFFER_SIZE` | @@ -507,7 +507,7 @@ | `cuModuleGetFunction` | `hipModuleGetFunction` | | `cuModuleGetGlobal` | `hipModuleGetGlobal` | | `cuModuleGetSurfRef` | | -| `cuModuleGetTexRef` | | +| `cuModuleGetTexRef` | `hipModuleGetTexRef` | | `cuModuleLoad` | `hipModuleLoad` | | `cuModuleLoadData` | `hipModuleLoadData` | | `cuModuleLoadDataEx` | `hipModuleLoadDataEx` | @@ -674,12 +674,12 @@ | `cuTexRefGetMipmappedArray` | | | `cuTexRefSetAddress` | | | `cuTexRefSetAddress2D` | | -| `cuTexRefSetAddressMode` | | -| `cuTexRefSetArray` | | +| `cuTexRefSetAddressMode` | `hipTexRefSetAddressMode` | +| `cuTexRefSetArray` | `hipTexRefSetArray` | | `cuTexRefSetBorderColor` | | -| `cuTexRefSetFilterMode` | | -| `cuTexRefSetFlags` | | -| `cuTexRefSetFormat` | | +| `cuTexRefSetFilterMode` | `hipTexRefSetFilterMode` | +| `cuTexRefSetFlags` | `hipTexRefSetFlags` | +| `cuTexRefSetFormat` | `hipTexRefSetFormat` | | `cuTexRefSetMaxAnisotropy` | | | `cuTexRefSetMipmapFilterMode` | | | `cuTexRefSetMipmapLevelBias` | | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..24c97258cf 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 75ed140298935b403b3493fa838fa7ba2e95751a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 18 Jan 2018 12:28:56 +0300 Subject: [PATCH 25/37] [HIPIFY] Add cudaMalloc3D support [ROCm/clr commit: eb27585136f832d335f0b0dc7b086e939f09e4d4] --- .../markdown/CUDA_Runtime_API_functions_supported_by_HIP.md | 2 +- projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index f27f1538b7..84fd0f469d 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -124,7 +124,7 @@ | `cudaHostRegister` | `hipHostRegister` | | `cudaHostUnregister` | `hipHostUnregister` | | `cudaMalloc` | `hipMalloc` | -| `cudaMalloc3D` | | +| `cudaMalloc3D` | `hipMalloc3D` | | `cudaMalloc3DArray` | `hipMalloc3DArray` | | `cudaMallocArray` | `hipMallocArray` | | `cudaMallocHost` | `hipHostMalloc` | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 4f468e7b47..b0c4d83706 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 fa4a5355e499c64913e37229637a8000c271ba37 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 19 Jan 2018 17:38:51 +0300 Subject: [PATCH 26/37] [HIPIFY] cuRAND lib support (partial - only Host) [ROCm/clr commit: 2e2e38f1639569f270b4ebf603d6a74e64dee9bf] --- .../markdown/CURAND_API_supported_by_HIP.md | 132 ++++++++++++++++++ .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 124 +++++++++++++++- .../hipamd/hipify-clang/src/Statistics.cpp | 4 +- .../clr/hipamd/hipify-clang/src/Statistics.h | 2 + 4 files changed, 259 insertions(+), 3 deletions(-) create mode 100644 projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md diff --git a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md new file mode 100644 index 0000000000..83c3d4a234 --- /dev/null +++ b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md @@ -0,0 +1,132 @@ +# CURAND API supported by HIP + +## **1. CURAND Data types** + +| **type** | **CUDA** | **HIP** | **HIP value** (if differs) | +|-------------:|---------------------------------------------------------------|------------------------------------------------------------|----------------------------| +| enum |***`curandStatus`*** |***`hiprandStatus`*** | +| enum |***`curandStatus_t`*** |***`hiprandStatus_t`*** | +| 0 |*`CURAND_STATUS_SUCCESS`* |*`HIPRAND_STATUS_SUCCESS`* | +| 100 |*`CURAND_STATUS_VERSION_MISMATCH`* |*`HIPRAND_STATUS_VERSION_MISMATCH`* | +| 101 |*`CURAND_STATUS_NOT_INITIALIZED`* |*`HIPRAND_STATUS_NOT_INITIALIZED`* | +| 102 |*`CURAND_STATUS_ALLOCATION_FAILED`* |*`HIPRAND_STATUS_ALLOCATION_FAILED`* | +| 103 |*`CURAND_STATUS_TYPE_ERROR`* |*`HIPRAND_STATUS_TYPE_ERROR`* | +| 104 |*`CURAND_STATUS_OUT_OF_RANGE`* |*`HIPRAND_STATUS_OUT_OF_RANGE`* | +| 105 |*`CURAND_STATUS_LENGTH_NOT_MULTIPLE`* |*`HIPRAND_STATUS_LENGTH_NOT_MULTIPLE`* | +| 106 |*`CURAND_STATUS_DOUBLE_PRECISION_REQUIRED`* |*`HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED`* | +| 201 |*`CURAND_STATUS_LAUNCH_FAILURE`* |*`HIPRAND_STATUS_LAUNCH_FAILURE`* | +| 202 |*`CURAND_STATUS_PREEXISTING_FAILURE`* |*`HIPRAND_STATUS_PREEXISTING_FAILURE`* | +| 203 |*`CURAND_STATUS_INITIALIZATION_FAILED`* |*`HIPRAND_STATUS_INITIALIZATION_FAILED`* | +| 204 |*`CURAND_STATUS_ARCH_MISMATCH`* |*`HIPRAND_STATUS_ARCH_MISMATCH`* | +| 999 |*`CURAND_STATUS_INTERNAL_ERROR`* |*`HIPRAND_STATUS_INTERNAL_ERROR`* | +| enum |***`curandRngType`*** |***`hiprandRngType`*** | +| enum |***`curandRngType_t`*** |***`hiprandRngType_t`*** | +| 0 |*`CURAND_RNG_TEST`* |*`HIPRAND_RNG_TEST`* | +| 100 |*`CURAND_RNG_PSEUDO_DEFAULT`* |*`HIPRAND_RNG_PSEUDO_DEFAULT`* | 400 | +| 101 |*`CURAND_RNG_PSEUDO_XORWOW`* |*`HIPRAND_RNG_PSEUDO_XORWOW`* | 401 | +| 121 |*`CURAND_RNG_PSEUDO_MRG32K3A`* |*`HIPRAND_RNG_PSEUDO_MRG32K3A`* | 402 | +| 141 |*`CURAND_RNG_PSEUDO_MTGP32`* |*`HIPRAND_RNG_PSEUDO_MTGP32`* | 403 | +| 142 |*`CURAND_RNG_PSEUDO_MT19937`* |*`HIPRAND_RNG_PSEUDO_MT19937`* | 404 | +| 161 |*`CURAND_RNG_PSEUDO_PHILOX4_32_10`* |*`HIPRAND_RNG_PSEUDO_PHILOX4_32_10`* | 405 | +| 200 |*`CURAND_RNG_QUASI_DEFAULT`* |*`HIPRAND_RNG_QUASI_DEFAULT`* | 500 | +| 201 |*`CURAND_RNG_QUASI_SOBOL32`* |*`HIPRAND_RNG_QUASI_SOBOL32`* | 501 | +| 202 |*`CURAND_RNG_QUASI_SCRAMBLED_SOBOL32`* |*`HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32`* | 502 | +| 203 |*`CURAND_RNG_QUASI_SOBOL64`* |*`HIPRAND_RNG_QUASI_SOBOL64`* | 503 | +| 204 |*`CURAND_RNG_QUASI_SCRAMBLED_SOBOL64`* |*`HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64`* | 504 | +| enum |***`curandOrdering`*** | | +| enum |***`curandOrdering_t`*** | | +| 100 |*`CURAND_ORDERING_PSEUDO_BEST`* | | +| 101 |*`CURAND_ORDERING_PSEUDO_DEFAULT`* | | +| 102 |*`CURAND_ORDERING_PSEUDO_SEEDED`* | | +| 201 |*`CURAND_ORDERING_QUASI_DEFAULT`* | | +| enum |***`curandDirectionVectorSet`*** | | +| enum |***`curandDirectionVectorSet_t`*** | | +| 101 |*`CURAND_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 102 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 103 |*`CURAND_DIRECTION_VECTORS_64_JOEKUO6`* | | +| 104 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`* | | +| uint | `curandDirectionVectors32_t` | `hiprandDirectionVectors32_t` | +| uint | `curandDirectionVectors64_t` | | +| struct | `curandGenerator_st` | `hiprandGenerator_st` | +| struct* | `curandGenerator_t` | `hiprandGenerator_t` | +| enum |***`curandOrdering`*** | | +| enum |***`curandOrdering_t`*** | | +| 100 |*`CURAND_ORDERING_PSEUDO_BEST`* | | +| 101 |*`CURAND_ORDERING_PSEUDO_DEFAULT`* | | +| 102 |*`CURAND_ORDERING_PSEUDO_SEEDED`* | | +| 201 |*`CURAND_ORDERING_QUASI_DEFAULT`* | | +| enum |***`curandDirectionVectorSet`*** | | +| enum |***`curandDirectionVectorSet_t`*** | | +| 101 |*`CURAND_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 102 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`* | | +| 103 |*`CURAND_DIRECTION_VECTORS_64_JOEKUO6`* | | +| 104 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`* | | +| uint | `curandDirectionVectors32_t` | `hiprandDirectionVectors32_t` | +| uint | `curandDirectionVectors64_t` | | +| double | `curandDistribution_st` | | +| double | `curandHistogramM2V_st` | | +| double* | `curandDistribution_t` | | +| double* | `curandHistogramM2V_t` | | +| struct | `curandDistributionShift_st` | | +| struct* | `curandDistributionShift_t` | | +| struct | `curandDistributionM2Shift_st` | | +| struct* | `curandDistributionM2Shift_t` | | +| struct | `curandHistogramM2_st` | | +| struct* | `curandHistogramM2_t` | | +| uint | `curandHistogramM2K_st` | | +| uint* | `curandHistogramM2K_t` | | +| struct | `curandDiscreteDistribution_st` | `hiprandDiscreteDistribution_st` | +| struct* | `curandDiscreteDistribution_t` | `hiprandDiscreteDistribution_t` | +| enum |***`curandMethod`*** | | +| enum |***`curandMethod_t`*** | | +| 0 |*`CURAND_CHOOSE_BEST`* | | +| 1 |*`CURAND_ITR`* | | +| 2 |*`CURAND_KNUTH`* | | +| 3 |*`CURAND_HITR`* | | +| 4 |*`CURAND_M1`* | | +| 5 |*`CURAND_M2`* | | +| 6 |*`CURAND_BINARY_SEARCH`* | | +| 7 |*`CURAND_DISCRETE_GAUSS`* | | +| 8 |*`CURAND_REJECTION`* | | +| 9 |*`CURAND_DEVICE_API`* | | +| 10 |*`CURAND_FAST_REJECTION`* | | +| 11 |*`CURAND_3RD`* | | +| 12 |*`CURAND_DEFINITION`* | | +| 13 |*`CURAND_POISSON`* | | + + +## **2. Host API Functions** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|--------------------------------------------| +| `curandCreateGenerator` | `hiprandCreateGenerator` | +| `curandCreateGeneratorHost` | `hiprandCreateGeneratorHost` | +| `curandCreatePoissonDistribution` | `hiprandCreatePoissonDistribution` | +| `curandDestroyDistribution` | `hiprandDestroyDistribution` | +| `curandDestroyGenerator` | `hiprandDestroyGenerator` | +| `curandGenerate` | `hiprandGenerate` | +| `curandGenerateLogNormal` | `hiprandGenerateLogNormal` | +| `curandGenerateLogNormalDouble` | `hiprandGenerateLogNormalDouble` | +| `curandGenerateLongLong` | | +| `curandGenerateNormal` | `hiprandGenerateNormal` | +| `curandGenerateNormalDouble` | `hiprandGenerateNormalDouble` | +| `curandGeneratePoisson` | `hiprandGeneratePoisson` | +| `curandGenerateSeeds` | `hiprandGenerateSeeds` | +| `curandGenerateUniform` | `hiprandGenerateUniform` | +| `curandGenerateUniformDouble` | `hiprandGenerateUniformDouble` | +| `curandGetDirectionVectors32` | | +| `curandGetDirectionVectors64` | | +| `curandGetProperty` | | +| `curandGetScrambleConstants32` | | +| `curandGetScrambleConstants64` | | +| `curandGetVersion` | `hiprandGetVersion` | +| `curandSetGeneratorOffset` | `hiprandSetGeneratorOffset` | +| `curandSetGeneratorOrdering` | | +| `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` | +| `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` | +| `curandSetStream` | `hiprandSetStream` | + +## **3. Device API Functions** + +| **CUDA** | **HIP** | +|-----------------------------------------------------------|-------------------------------| diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 755ce365af..b410ec43f9 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -323,6 +323,36 @@ const std::map CUDA_TYPE_NAME_MAP{ {"cublasPointerMode_t", {"hipblasPointerMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasAtomicsMode_t", {"hipblasAtomicsMode_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, {"cublasDataType_t", {"hipblasDataType_t", CONV_TYPE, API_BLAS, HIP_UNSUPPORTED}}, + + ///////////////////////////// cuRAND ///////////////////////////// + {"curandStatus", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, + {"curandStatus_t", {"hiprandStatus_t", CONV_TYPE, API_RAND}}, + {"curandRngType", {"hiprandRngType_t", CONV_TYPE, API_RAND}}, + {"curandRngType_t", {"hiprandRngType_t", CONV_TYPE, API_RAND}}, + {"curandGenerator_st", {"hiprandGenerator_st", CONV_TYPE, API_RAND}}, + {"curandGenerator_t", {"hiprandGenerator_t", CONV_TYPE, API_RAND}}, + {"curandDirectionVectorSet", {"hiprandDirectionVectorSet_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDirectionVectorSet_t", {"hiprandDirectionVectorSet_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandOrdering", {"hiprandOrdering_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandOrdering_t", {"hiprandOrdering_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistribution_st", {"hiprandDistribution_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2V_st", {"hiprandDistribution_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistribution_t", {"hiprandDistribution_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2V_t", {"hiprandDistribution_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionShift_st", {"hiprandDistributionShift_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionShift_t", {"hiprandDistributionShift_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionM2Shift_st", {"hiprandDistributionM2Shift_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDistributionM2Shift_t", {"hiprandDistributionM2Shift_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2_st", {"hiprandHistogramM2_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2_t", {"hiprandHistogramM2_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2K_st", {"hiprandHistogramM2K_st", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandHistogramM2K_t", {"hiprandHistogramM2K_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDiscreteDistribution_st", {"hiprandDiscreteDistribution_st", CONV_TYPE, API_RAND}}, + {"curandDiscreteDistribution_t", {"hiprandDiscreteDistribution_t", CONV_TYPE, API_RAND}}, + {"curandMethod", {"hiprandMethod_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandMethod_t", {"hiprandMethod_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, + {"curandDirectionVectors32_t", {"hiprandDirectionVectors32_t", CONV_TYPE, API_RAND}}, + {"curandDirectionVectors64_t", {"hiprandDirectionVectors64_t", CONV_TYPE, API_RAND, HIP_UNSUPPORTED}}, }; /// Maps cuda header names to hip header names. @@ -343,6 +373,10 @@ const std::map CUDA_INCLUDE_MAP{ {"cublas.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, + // CURAND includes + {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, + {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + // HIP includes // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h // {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}}, @@ -2660,7 +2694,95 @@ const std::map CUDA_IDENTIFIER_MAP{ // ROTMG {"cublasSrotmg_v2", {"hipblasSrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, - {"cublasDrotmg_v2", {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}} + {"cublasDrotmg_v2", {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS, HIP_UNSUPPORTED}}, + + ///////////////////////////// cuRAND ///////////////////////////// + // RAND function call status types (enum curandStatus) + {"CURAND_STATUS_SUCCESS", {"HIPRAND_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_VERSION_MISMATCH", {"HIPRAND_STATUS_VERSION_MISMATCH", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_NOT_INITIALIZED", {"HIPRAND_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_ALLOCATION_FAILED", {"HIPRAND_STATUS_ALLOCATION_FAILED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_TYPE_ERROR", {"HIPRAND_STATUS_TYPE_ERROR", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_OUT_OF_RANGE", {"HIPRAND_STATUS_OUT_OF_RANGE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_LENGTH_NOT_MULTIPLE", {"HIPRAND_STATUS_LENGTH_NOT_MULTIPLE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_DOUBLE_PRECISION_REQUIRED", {"HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_LAUNCH_FAILURE", {"HIPRAND_STATUS_LAUNCH_FAILURE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_PREEXISTING_FAILURE", {"HIPRAND_STATUS_PREEXISTING_FAILURE", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_INITIALIZATION_FAILED", {"HIPRAND_STATUS_INITIALIZATION_FAILED", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_ARCH_MISMATCH", {"HIPRAND_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_STATUS_INTERNAL_ERROR", {"HIPRAND_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_RAND}}, + + // RAND generator types (enum curandRngType) + {"CURAND_RNG_TEST", {"HIPRAND_RNG_TEST", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_DEFAULT", {"HIPRAND_RNG_PSEUDO_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_XORWOW", {"HIPRAND_RNG_PSEUDO_XORWOW", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MRG32K3A", {"HIPRAND_RNG_PSEUDO_MRG32K3A", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MTGP32", {"HIPRAND_RNG_PSEUDO_MTGP32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_MT19937", {"HIPRAND_RNG_PSEUDO_MT19937", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_PSEUDO_PHILOX4_32_10", {"HIPRAND_RNG_PSEUDO_PHILOX4_32_10", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_DEFAULT", {"HIPRAND_RNG_QUASI_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SOBOL32", {"HIPRAND_RNG_QUASI_SOBOL32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SCRAMBLED_SOBOL32", {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SOBOL64", {"HIPRAND_RNG_QUASI_SOBOL64", CONV_NUMERIC_LITERAL, API_RAND}}, + {"CURAND_RNG_QUASI_SCRAMBLED_SOBOL64", {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64", CONV_NUMERIC_LITERAL, API_RAND}}, + + // RAND ordering of results in memory (enum curandOrdering) + {"CURAND_ORDERING_PSEUDO_BEST", {"HIPRAND_ORDERING_PSEUDO_BEST", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_PSEUDO_DEFAULT", {"HIPRAND_ORDERING_PSEUDO_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_PSEUDO_SEEDED", {"HIPRAND_ORDERING_PSEUDO_SEEDED", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ORDERING_QUASI_DEFAULT", {"HIPRAND_ORDERING_QUASI_DEFAULT", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND choice of direction vector set (enum curandDirectionVectorSet) + {"CURAND_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND method (enum curandMethod) + {"CURAND_CHOOSE_BEST", {"HIPRAND_CHOOSE_BEST", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_ITR", {"HIPRAND_ITR", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_KNUTH", {"HIPRAND_KNUTH", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_HITR", {"HIPRAND_HITR", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_M1", {"HIPRAND_M1", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_M2", {"HIPRAND_M2", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_BINARY_SEARCH", {"HIPRAND_BINARY_SEARCH", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DISCRETE_GAUSS", {"HIPRAND_DISCRETE_GAUSS", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_REJECTION", {"HIPRAND_REJECTION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DEVICE_API", {"HIPRAND_DEVICE_API", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_FAST_REJECTION", {"HIPRAND_FAST_REJECTION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_3RD", {"HIPRAND_3RD", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_DEFINITION", {"HIPRAND_DEFINITION", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + {"CURAND_POISSON", {"HIPRAND_POISSON", CONV_NUMERIC_LITERAL, API_RAND, HIP_UNSUPPORTED}}, + + // RAND Host functions + {"curandCreateGenerator", {"hiprandCreateGenerator", CONV_MATH_FUNC, API_RAND}}, + {"curandCreateGeneratorHost", {"hiprandCreateGeneratorHost", CONV_MATH_FUNC, API_RAND}}, + {"curandCreatePoissonDistribution", {"hiprandCreatePoissonDistribution", CONV_MATH_FUNC, API_RAND}}, + {"curandDestroyDistribution", {"hiprandDestroyDistribution", CONV_MATH_FUNC, API_RAND}}, + {"curandDestroyGenerator", {"hiprandDestroyGenerator", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerate", {"hiprandGenerate", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLogNormal", {"hiprandGenerateLogNormal", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLogNormalDouble", {"hiprandGenerateLogNormalDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateLongLong", {"hiprandGenerateLongLong", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGenerateNormal", {"hiprandGenerateNormal", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateNormalDouble", {"hiprandGenerateNormalDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGeneratePoisson", {"hiprandGeneratePoisson", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateSeeds", {"hiprandGenerateSeeds", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateUniform", {"hiprandGenerateUniform", CONV_MATH_FUNC, API_RAND}}, + {"curandGenerateUniformDouble", {"hiprandGenerateUniformDouble", CONV_MATH_FUNC, API_RAND}}, + {"curandGetDirectionVectors32", {"hiprandGetDirectionVectors32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetDirectionVectors64", {"hiprandGetDirectionVectors64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetProperty", {"hiprandGetProperty", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetScrambleConstants32", {"hiprandGetScrambleConstants32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetScrambleConstants64", {"hiprandGetScrambleConstants64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandGetVersion", {"hiprandGetVersion", CONV_MATH_FUNC, API_RAND}}, + {"curandSetGeneratorOffset", {"hiprandSetGeneratorOffset", CONV_MATH_FUNC, API_RAND}}, + {"curandSetGeneratorOrdering", {"hiprandSetGeneratorOrdering", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}}, + {"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", CONV_MATH_FUNC, API_RAND}}, + {"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", CONV_MATH_FUNC, API_RAND}}, + {"curandSetStream", {"hiprandSetStream", CONV_MATH_FUNC, API_RAND}}, + + // RAND Device functions }; const std::map& CUDA_RENAMES_MAP() { diff --git a/projects/clr/hipamd/hipify-clang/src/Statistics.cpp b/projects/clr/hipamd/hipify-clang/src/Statistics.cpp index 00b8a66f11..e2b3632066 100644 --- a/projects/clr/hipamd/hipify-clang/src/Statistics.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Statistics.cpp @@ -5,7 +5,7 @@ const char *counterNames[NUM_CONV_TYPES] = { - "version", "init", "device", "mem", "kern", "coord_func", "math_func", + "version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func", "special_func", "stream", "event", "occupancy", "ctx", "peer", "module", "cache", "exec", "err", "def", "tex", "gl", "graphics", "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", @@ -14,7 +14,7 @@ const char *counterNames[NUM_CONV_TYPES] = { }; const char *apiNames[NUM_API_TYPES] = { - "CUDA Driver API", "CUDA RT API", "CUBLAS API" + "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API" }; namespace { diff --git a/projects/clr/hipamd/hipify-clang/src/Statistics.h b/projects/clr/hipamd/hipify-clang/src/Statistics.h index da4e296db0..f160ca7383 100644 --- a/projects/clr/hipamd/hipify-clang/src/Statistics.h +++ b/projects/clr/hipamd/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 51a5587d2503e72b4903eb952d7f2308123f03d2 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 19 Jan 2018 21:29:05 +0300 Subject: [PATCH 27/37] [HIPIFY] cuRAND lib support (Device) [ROCm/clr commit: 5b1b3438fd072edb68ae6ed6d911fb743476ed07] --- .../markdown/CURAND_API_supported_by_HIP.md | 41 +++++++++++++++++-- .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 37 ++++++++++++++++- 2 files changed, 74 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md index 83c3d4a234..900cfadc03 100644 --- a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md @@ -93,7 +93,14 @@ | 11 |*`CURAND_3RD`* | | | 12 |*`CURAND_DEFINITION`* | | | 13 |*`CURAND_POISSON`* | | - +| struct | `curandStateMtgp32_t` | `hiprandStateMtgp32_t` | +| struct | `curandStateScrambledSobol64_t` | | +| struct | `curandStateSobol64_t` | | +| struct | `curandStateScrambledSobol32_t` | | +| struct | `curandStateSobol32_t` | `hiprandStateSobol32_t` | +| struct | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` | +| struct | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` | +| struct | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` | ## **2. Host API Functions** @@ -128,5 +135,33 @@ ## **3. Device API Functions** -| **CUDA** | **HIP** | -|-----------------------------------------------------------|-------------------------------| +| **CUDA** | **HIP** | +|-----------------------------------------------------------|--------------------------------------------| +| `curand` | `hiprand` | +| `curand_init` | `hiprand_init` | +| `curand_log_normal` | `hiprand_log_normal` | +| `curand_log_normal_double` | `hiprand_log_normal_double` | +| `curand_log_normal2` | `hiprand_log_normal2` | +| `curand_log_normal2_double` | `hiprand_log_normal2_double` | +| `curand_log_normal4` | `hiprand_log_normal4` | +| `curand_log_normal4_double` | `hiprand_log_normal4_double` | +| `curand_mtgp32_single` | | +| `curand_mtgp32_single_specific` | | +| `curand_mtgp32_specific` | | +| `curand_normal` | `hiprand_normal` | +| `curand_normal_double` | `hiprand_normal_double` | +| `curand_normal2` | `hiprand_normal2` | +| `curand_normal2_double` | `hiprand_normal2_double` | +| `curand_normal4` | `hiprand_normal4` | +| `curand_normal4_double` | `hiprand_normal4_double` | +| `curand_uniform2_double` | `hiprand_uniform2_double` | +| `curand_uniform4` | `hiprand_uniform4` | +| `curand_uniform4_double` | `hiprand_uniform4_double` | +| `curand_discrete` | `hiprand_discrete` | +| `curand_discrete4` | `hiprand_discrete4` | +| `curand_poisson` | `hiprand_poisson` | +| `curand_poisson4` | `hiprand_poisson4` | +| `curand_Philox4x32_10` | | +| `skipahead` | `skipahead` | +| `skipahead_sequence` | `skipahead_sequence` | +| `skipahead_subsequence` | `skipahead_subsequence` | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index b410ec43f9..db78660121 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 d88edc88740652a2dc97af0eddfac8b1bbe73597 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 22 Jan 2018 17:12:02 +0300 Subject: [PATCH 28/37] [HIPIFY] Sync with hipBLAS ToT and CUDA cuBLAS 9.1 [ROCm/clr commit: 1ecb5b5901e5a970ae0cff24a8905e78e61fa2b9] --- .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 97 +++++++++++++------ 1 file changed, 70 insertions(+), 27 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index db78660121..1f9c6287ed 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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 34b797d7f238f6950a5f8efddb1cd193fc71b675 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 21:43:18 +0300 Subject: [PATCH 29/37] [HIPIFY][tests][win] Fix run_test.bat All checks should not occur in input file for FileCheck. The issue found on CHECK-NOT. Change removes all lit checks in the hipified file based on regexp, and the resulted stdout is fed as stdin for FileCheck. [ROCm/clr commit: 8fae5663d279bfe420aa8968a43b2f812b5bda7a] --- projects/clr/hipamd/tests/hipify-clang/run_test.bat | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.bat b/projects/clr/hipamd/tests/hipify-clang/run_test.bat index 6eefb7e46e..b4858db807 100644 --- a/projects/clr/hipamd/tests/hipify-clang/run_test.bat +++ b/projects/clr/hipamd/tests/hipify-clang/run_test.bat @@ -1,4 +1,4 @@ -@echo off +rem @echo off setlocal for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i @@ -14,5 +14,6 @@ set clang_args=%4%clang_args% %HIPIFY% -o=%TMP_FILE% %IN_FILE% -- %clang_args% if errorlevel 1 (echo Error: hipify-clang.exe failed with exit code: %errorlevel% && exit /b %errorlevel%) -%FILE_CHECK% %IN_FILE% -input-file=%TMP_FILE% + +findstr /v /r /c:"[ ]*//[ ]*[CHECK*|RUN]" %TMP_FILE% | %FILE_CHECK% %IN_FILE% if errorlevel 1 (echo Error: FileCheck.exe failed with exit code: %errorlevel% && exit /b %errorlevel%) From f78369ce58e13bd33642f692a6995ebef7e37d2e Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 21:46:27 +0300 Subject: [PATCH 30/37] [HIPIFY][tests][win] Uncomment @echo off [ROCm/clr commit: 35bd23fb07a1bd48f09749473827de782aa6eda1] --- projects/clr/hipamd/tests/hipify-clang/run_test.bat | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.bat b/projects/clr/hipamd/tests/hipify-clang/run_test.bat index b4858db807..d8c8d74cf0 100644 --- a/projects/clr/hipamd/tests/hipify-clang/run_test.bat +++ b/projects/clr/hipamd/tests/hipify-clang/run_test.bat @@ -1,4 +1,4 @@ -rem @echo off +@echo off setlocal for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i From d3d192e62cf8e2a21e117279ab009d04a3a2b2a5 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 23:06:55 +0300 Subject: [PATCH 31/37] [HIPIFY][fix] Fix PragmaDirective File location have to be verified, otherwise location of the first found '#pragma once' in any included header even system will be erroneously handled, which might lead to attempt to including hip_runtime.h in it. [ROCm/clr commit: b9fc653b089e0827cc2f95506bd165e67abe228b] --- projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..7d9ddebf51 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -198,8 +198,13 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, } void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) { - if (pragmaOnce) { return; } + if (pragmaOnce) { + return; + } clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(Loc)) { + return; + } clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); const clang::Token tok = PP.LookAhead(0); StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); From cbbf54d1221b05c98f55bd5e560658a3de16a1a1 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 23:43:36 +0300 Subject: [PATCH 32/37] [HIPIFY][fix] CUDA and cuBLAS main headers correct handling [ROCm/clr commit: aaa95597689700a0c0117a7e8fb18e4e7987b600] --- .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 8 ++--- .../hipamd/hipify-clang/src/HipifyAction.cpp | 35 +++++++++++++------ .../hipamd/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 projects/clr/hipamd/tests/hipify-clang/headers_test_06.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/headers_test_07.cu create mode 100644 projects/clr/hipamd/tests/hipify-clang/headers_test_08.cu diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 1f9c6287ed..6bcc0b38c4 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -379,12 +379,12 @@ const std::map CUDA_INCLUDE_MAP{ {"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}}, // CUBLAS includes - {"cublas.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, - {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, + {"cublas.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, // CURAND includes - {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, - {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, + {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, // HIP includes // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..9ac0eacbd1 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -149,23 +149,38 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, if (!SM.isWrittenInMainFile(hash_loc)) { return; } + if (!firstHeader) { + firstHeader = true; + firstHeaderLoc = hash_loc; + } const auto found = CUDA_INCLUDE_MAP.find(file_name); if (found == CUDA_INCLUDE_MAP.end()) { - if (!firstNotMainHeader) { - firstNotMainHeader = true; - firstNotMainHeaderLoc = hash_loc; - } return; } // Special-casing to avoid duplication of the hip_runtime include. bool secondMainInclude = false; - if (found->second.hipName == "hip/hip_runtime.h") { - if (insertedRuntimeHeader) { - secondMainInclude = true; + if (found->second.countType == CONV_INCLUDE_CUDA_MAIN_H) { + switch (found->second.countApiType) { + case API_DRIVER: + case API_RUNTIME: + if (insertedRuntimeHeader) { + secondMainInclude = true; + break; + } + insertedRuntimeHeader = true; + break; + case API_BLAS: + if (insertedBLASHeader) { + secondMainInclude = true; + break; + } + insertedBLASHeader = true; + break; + default: + break; } - insertedRuntimeHeader = true; } Statistics::current().incrementCounter(found->second, file_name.str()); @@ -356,8 +371,8 @@ void HipifyAction::EndSourceFileAction() { clang::SourceLocation sl; if (pragmaOnce) { sl = pragmaOnceLoc; - } else if (firstNotMainHeader) { - sl = firstNotMainHeaderLoc; + } else if (firstHeader) { + sl = firstHeaderLoc; } else { sl = SM.getLocForStartOfFile(SM.getMainFileID()); } diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h index a269a37117..42622c1e01 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h @@ -23,9 +23,10 @@ private: // not, we insert it at the top of the file when we finish processing it. // This approach means we do the best it's possible to do w.r.t preserving the user's include order. bool insertedRuntimeHeader = false; - bool firstNotMainHeader = false; + bool insertedBLASHeader = false; + bool firstHeader = false; bool pragmaOnce = false; - clang::SourceLocation firstNotMainHeaderLoc; + clang::SourceLocation firstHeaderLoc; clang::SourceLocation pragmaOnceLoc; /** diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_01.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_01.cu index c39ef80d8f..3747c339e8 100644 --- a/projects/clr/hipamd/tests/hipify-clang/headers_test_01.cu +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_01.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args // CHECK: #include +// CHECK-NOT: #include +// CHECK: #include #include -// CHECK-NOT: #include #include +#include diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_02.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_02.cu index 90d412f797..57308efd59 100644 --- a/projects/clr/hipamd/tests/hipify-clang/headers_test_02.cu +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_02.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args -// CHECK: #include -#include -// CHECK-NOT: #include -#include +// CHECK: #include "hip/hip_runtime.h" +// CHECK-NOT: #include "cuda_runtime.h" +// CHECK: #include +#include "cuda.h" +#include "cuda_runtime.h" +#include diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_06.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_06.cu new file mode 100644 index 0000000000..bce73c42df --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_06.cu @@ -0,0 +1,8 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +#include +#include +#include diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_07.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_07.cu new file mode 100644 index 0000000000..4237e1eb72 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_07.cu @@ -0,0 +1,8 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" +// CHECK: #include +#include "cublas_v2.h" +#include "cublas.h" +#include diff --git a/projects/clr/hipamd/tests/hipify-clang/headers_test_08.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_08.cu new file mode 100644 index 0000000000..ad54871bd8 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/headers_test_08.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" +// CHECK: #include +#include +#include +#include +#include "cublas_v2.h" +#include "cublas.h" +#include From f33c431dd985e9f731d589c861a4566ea4ca1d4b Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 24 Jan 2018 20:13:23 +0300 Subject: [PATCH 33/37] [HIPIFY][tests][win] Make cudaRegister.cu building on Windows as well [ROCm/clr commit: 4df35f4c778365d72680bf7615807afa92120a14] --- projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu b/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu index 79d21707c2..43b4345337 100644 --- a/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu +++ b/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu @@ -22,7 +22,12 @@ THE SOFTWARE. #include #include #include -#include +#ifdef _WIN32 +#include +#define sleep(x) Sleep(x) +#else +#include +#endif #include #include @@ -33,7 +38,7 @@ THE SOFTWARE. // CHECK: if(status != hipSuccess) { #define check(msg, status){ \ if(status != cudaSuccess) { \ - printf("%s failed. \n", #msg); \ + printf("%s failed. \n", #msg); \ } \ } From 7df2ff97aeea93d8afcb969133df92e94aea2216 Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Wed, 24 Jan 2018 17:00:57 -0600 Subject: [PATCH 34/37] Fixing rocblas build failure with ::Bundled_code_header constructor Disabling hipPrintfKernel test from CI [ROCm/clr commit: 85284dd48fcf5311db775e1dc752ce0389bd0446] --- projects/clr/hipamd/Jenkinsfile | 3 ++- .../clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/Jenkinsfile b/projects/clr/hipamd/Jenkinsfile index 4909e666f0..6142b94d95 100644 --- a/projects/clr/hipamd/Jenkinsfile +++ b/projects/clr/hipamd/Jenkinsfile @@ -160,6 +160,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf } // Cap the maximum amount of testing, in case of hangs + // Excluding hipPrintfKernel test from automation; variable fails on CI test machines timeout(time: 1, unit: 'HOURS') { stage("${platform} unit testing") @@ -169,7 +170,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - make test + ctest -E hipPrintfKernel """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard diff --git a/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp b/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..2bec0017db 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp @@ -150,8 +150,8 @@ namespace hip_impl }; // CREATORS - template - Bundled_code_header::Bundled_code_header(I f, I l) : Bundled_code_header{} + template + Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l) : Bundled_code_header{} { read(f, l, *this); } From 98595dcc9d1f4046ab6ae6bb4c635b9ac22beb3d Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 25 Jan 2018 16:51:29 +0000 Subject: [PATCH 35/37] use assign rather than insert [ROCm/clr commit: 6f90e4213f10dae5baeb24b6dd7b28e5564aa448] --- .../clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp b/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp index 72f9d35c73..1b2dfc1c7c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/code_object_bundle.hpp @@ -88,7 +88,7 @@ namespace hip_impl std::copy_n(it, sizeof(y.cbuf), y.cbuf); it += sizeof(y.cbuf); - y.triple.insert(y.triple.cend(), it, it + y.triple_sz); + y.triple.assign(it, it + y.triple_sz); std::copy_n( f + y.offset, y.bundle_sz, std::back_inserter(y.blob)); @@ -155,4 +155,4 @@ namespace hip_impl { read(f, l, *this); } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. From 978d483eb5357aa66b6201a99dd9cadf3cacfc96 Mon Sep 17 00:00:00 2001 From: Kent Knox Date: Thu, 25 Jan 2018 10:52:56 -0600 Subject: [PATCH 36/37] Remove archiving of RPM We should archive RPM's from proper centos/fedora machines so that we get the proper dependencies right [ROCm/clr commit: 2c89a70d3f0ea7992d2774d0142c62102de13af0] --- projects/clr/hipamd/Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/Jenkinsfile b/projects/clr/hipamd/Jenkinsfile index 6142b94d95..6e5f7bc8e5 100644 --- a/projects/clr/hipamd/Jenkinsfile +++ b/projects/clr/hipamd/Jenkinsfile @@ -151,7 +151,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf // The rm command needs to run as sudo because the test steps below create files owned by root sh """#!/usr/bin/env bash set -x - sudo rm -rf ${build_dir_rel} + rm -rf ${build_dir_rel} mkdir -p ${build_dir_rel} cd ${build_dir_rel} cmake -DCMAKE_BUILD_TYPE=${build_config} -DCMAKE_INSTALL_PREFIX=staging ${optional_configure} ${source_hip_abs} @@ -194,7 +194,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf if( platform.toLowerCase( ).startsWith( 'hcc-ctu' ) ) { archiveArtifacts artifacts: "${build_dir_rel}/*.deb", fingerprint: true - archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true + // archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true } } } From d4ff298f5e7343e8055f5f169ca3a7ae7c7bcd14 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 26 Jan 2018 07:02:49 +0530 Subject: [PATCH 37/37] Disable md2html conversion in hip doc package [ROCm/clr commit: 4995146eaff8358e25c2c7673a4a98266bd92a23] --- projects/clr/hipamd/packaging/hip_doc.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/packaging/hip_doc.txt b/projects/clr/hipamd/packaging/hip_doc.txt index daef7810b4..5987c2c128 100644 --- a/projects/clr/hipamd/packaging/hip_doc.txt +++ b/projects/clr/hipamd/packaging/hip_doc.txt @@ -8,12 +8,12 @@ if(DOXYGEN_EXE) install(DIRECTORY RuntimeAPI/html DESTINATION docs/docs/RuntimeAPI) endif() -find_program(GRIP_EXE grip) -if(GRIP_EXE) - add_custom_target(convert_md_to_html ALL - COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html) - install(DIRECTORY md2html/ DESTINATION docs) -endif() +#find_program(GRIP_EXE grip) +#if(GRIP_EXE) +# add_custom_target(convert_md_to_html ALL +# COMMAND @hip_SOURCE_DIR@/packaging/convert_md_to_html.sh @hip_SOURCE_DIR@ ${PROJECT_BINARY_DIR}/md2html) +# install(DIRECTORY md2html/ DESTINATION docs) +#endif() ############################# # Packaging steps