diff --git a/CMakeLists.txt b/CMakeLists.txt index 8278daad2c..c62ec8d85d 100755 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -33,30 +33,57 @@ string(REPLACE "." ";" VERSION_LIST ${HIP_BASE_VERSION}) list(GET VERSION_LIST 0 HIP_VERSION_MAJOR) list(GET VERSION_LIST 1 HIP_VERSION_MINOR) -# get date information based on UTC -# use the last two digits of year + week number + day in the week as HIP_VERSION_GITDATE -# use the commit date, instead of build date -# add xargs to remove strange trailing newline character -execute_process(COMMAND git show -s --format=@%ct +find_package(Git) + +# FIXME: Two different version strings used. +if(GIT_FOUND) + # get date information based on UTC + # use the last two digits of year + week number + day in the week as HIP_VERSION_GITDATE + # use the commit date, instead of build date + # add xargs to remove strange trailing newline character + execute_process(COMMAND ${GIT_EXECUTABLE} show -s --format=@%ct COMMAND xargs COMMAND date -f - --utc +%y%U%w + RESULT_VARIABLE git_result + OUTPUT_VARIABLE git_output WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} - OUTPUT_VARIABLE HIP_VERSION_GITDATE OUTPUT_STRIP_TRAILING_WHITESPACE) + if(git_result EQUAL 0) + set(HIP_VERSION_GITDATE ${git_output}) + endif() -# get commit short hash -execute_process(COMMAND git rev-parse --short HEAD + # get commit short hash + execute_process(COMMAND ${GIT_EXECUTABLE} rev-parse --short HEAD WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} - OUTPUT_VARIABLE HIP_VERSION_GITHASH + RESULT_VARIABLE git_result + OUTPUT_VARIABLE git_output OUTPUT_STRIP_TRAILING_WHITESPACE) + if(git_result EQUAL 0) + set(HIP_VERSION_GITHASH ${git_output}) + endif() -# get commit count -execute_process(COMMAND git rev-list --count HEAD + # get commit count + execute_process(COMMAND ${GIT_EXECUTABLE} rev-list --count HEAD WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} - OUTPUT_VARIABLE HIP_VERSION_GITCOUNT + RESULT_VARIABLE git_result + OUTPUT_VARIABLE git_output OUTPUT_STRIP_TRAILING_WHITESPACE) + if(git_result EQUAL 0) + set(HIP_VERSION_GITCOUNT ${git_output}) + endif() + + set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}-${HIP_VERSION_GITHASH}) + + if(DEFINED ENV{ROCM_BUILD_ID}) + set(HIP_PACKAGING_VERSION_PATCH ${HIP_VERSION_GITDATE}.${HIP_VERSION_GITCOUNT}-$ENV{ROCM_BUILD_ID}-${HIP_VERSION_GITHASH}) + else() + set(HIP_PACKAGING_VERSION_PATCH ${HIP_VERSION_GITDATE}.${HIP_VERSION_GITCOUNT}-${HIP_VERSION_GITHASH}) + endif() +else() + # FIXME: Some parts depend on this being set. + set(HIP_PACKAGING_VERSION_PATCH "0") +endif() -set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}-${HIP_VERSION_GITHASH}) add_to_config(_versionInfo HIP_VERSION_MAJOR) add_to_config(_versionInfo HIP_VERSION_MINOR) add_to_config(_versionInfo HIP_VERSION_PATCH) @@ -85,12 +112,7 @@ if(CMAKE_CXX_COMPILER MATCHES ".*hcc") endif() # overwrite HIP_VERSION_PATCH for packaging -if(DEFINED ENV{ROCM_BUILD_ID}) - set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}.${HIP_VERSION_GITCOUNT}-$ENV{ROCM_BUILD_ID}-${HIP_VERSION_GITHASH}) -else() - set(HIP_VERSION_PATCH ${HIP_VERSION_GITDATE}.${HIP_VERSION_GITCOUNT}-${HIP_VERSION_GITHASH}) -endif() -set(HIP_VERSION ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_VERSION_PATCH}) +set(HIP_VERSION ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR}.${HIP_PACKAGING_VERSION_PATCH}) ############################# # Configure variables @@ -415,6 +437,10 @@ endif() # Generate .hipVersion file(WRITE "${PROJECT_BINARY_DIR}/.hipVersion" ${_versionInfo}) +if(NOT DEFINED HIP_VERSION_GITDATE) + set(HIP_VERSION_GITDATE 0) +endif() + # Generate hip_version.h set(_versionInfoHeader "// Auto-generated by cmake\n diff --git a/cmake/FindHIP/run_hipcc.cmake b/cmake/FindHIP/run_hipcc.cmake index 1d8f3b38ad..63c7bb3887 100644 --- a/cmake/FindHIP/run_hipcc.cmake +++ b/cmake/FindHIP/run_hipcc.cmake @@ -58,7 +58,8 @@ if(NOT host_flag) if(NOT "x${HIP_CLANG_PATH}" STREQUAL "x") set(ENV{HIP_CLANG_PATH} ${HIP_CLANG_PATH}) endif() - set(__CC_FLAGS ${HIP_CLANG_PARALLEL_BUILD_COMPILE_OPTIONS} ${HIP_HIPCC_FLAGS} ${HIP_CLANG_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_CLANG_FLAGS_${build_configuration}}) + # Temporarily include HIP_HCC_FLAGS for HIP-Clang for PyTorch builds + set(__CC_FLAGS ${HIP_CLANG_PARALLEL_BUILD_COMPILE_OPTIONS} ${HIP_HIPCC_FLAGS} ${HIP_HCC_FLAGS} ${HIP_CLANG_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_HCC_FLAGS_${build_configuration}} ${HIP_CLANG_FLAGS_${build_configuration}}) endif() else() set(__CC_FLAGS ${HIP_HIPCC_FLAGS} ${HIP_NVCC_FLAGS} ${HIP_HIPCC_FLAGS_${build_configuration}} ${HIP_NVCC_FLAGS_${build_configuration}}) diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index f8a7489fb7..8aa3daa239 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -15,20 +15,23 @@ - [Do HIPIFY tools automatically convert all source code?](#do-hipify-tools-automatically-convert-all-source-code) - [What is NVCC?](#what-is-nvcc) - [What is HCC?](#what-is-hcc) +- [What is HIP-Clang?](#what-is-hip-clang) - [Why use HIP rather than supporting CUDA directly?](#why-use-hip-rather-than-supporting-cuda-directly) - [Can I develop HIP code on an Nvidia CUDA platform?](#can-i-develop-hip-code-on-an-nvidia-cuda-platform) -- [Can I develop HIP code on an AMD HCC platform?](#can-i-develop-hip-code-on-an-amd-hcc-platform) +- [Can I develop HIP code on an AMD HIP-Clang platform?](#can-i-develop-hip-code-on-an-amd-hip-clang-platform) +- [Do I need to make code changes in HIP code if switch compiler from HCC to HIP-Clang?](#Do-I-need-to-make-code-changes-in-hip-code-if-switch-compiler-from-hcc-to-hip-clang) +- [How to use HIP-Clang to build HIP programs instead of HCC?](#how-to-use-hip-clang-to-build-hip-programs-instead-of-hcc) +- [What is ROCclr?](#what-is-rocclr) - [Can a HIP binary run on both AMD and Nvidia platforms?](#can-a-hip-binary-run-on-both-amd-and-nvidia-platforms) - [What's the difference between HIP and hc?](#whats-the-difference-between-hip-and-hc) -- [On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang?](#on-hcc-can-i-link-hip-code-with-host-code-compiled-with-another-compiler-such-as-gcc-icc-or-clang-) -- [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do) -- [Can I install both CUDA SDK and HCC on same machine?](#can-i-install-both-cuda-sdk-and-hcc-on-same-machine) +- [On HIP-Clang, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang?](#on-HIP-Clang-can-i-link-hip-code-with-host-code-compiled-with-another-compiler-such-as-gcc-icc-or-clang-) +- [HIP detected my platform (hip-clang vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hip-clang-vs-nvcc-incorrectly---what-should-i-do) +- [Can I install both CUDA SDK and HIP-clang on same machine?](#can-i-install-both-cuda-sdk-and-hip-clang-on-same-machine) - [On CUDA, can I mix CUDA code with HIP code?](#on-cuda-can-i-mix-cuda-code-with-hip-code) -- [On HCC, can I use HC functionality with HIP?](#on-hcc-can-i-use-hc-functionality-with-hip) +- [On HIP-Clang, can I use HC functionality with HIP?](#on-hip-clang-can-i-use-hc-functionality-with-hip) - [How do I trace HIP application flow?](#how-do-i-trace-hip-application-flow) - [What if HIP generates an error of "symbol multiply defined!" only on AMD machine?](#what-if-hip-generates-error-of-symbol-multiply-defined-only-on-amd-machine) -- [How do I disable HIP Generic Grid Launch option?](#how-do-i-disable-hip-generic-grid-launch-option) - +- [What is maximum limit of Generic kernel launching parameter?](#what-is-maximum-limit-of-generic-kernel-launching-parameter) ### What APIs and features does HIP support? @@ -47,6 +50,7 @@ HIP provides the following: The HIP API documentation describes each API and its limitations, if any, compared with the equivalent CUDA API. ### What is not supported? + #### Runtime/Driver API features At a high-level, the following features are not supported: - Textures (partial support available) @@ -62,7 +66,7 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for #### Kernel language features - C++-style device-side dynamic memory allocations (free, new, delete) (CUDA 4.0) - Virtual functions, indirect functions and try/catch (CUDA 4.0) -- `__prof_trigger` +- `__prof_trigger` - PTX assembly (CUDA 4.0). HCC supports inline GCN assembly. - Several kernel features are under development. See the [HIP Kernel Language](hip_kernel_language.md) for more information. These include: - printf @@ -70,19 +74,18 @@ See the [API Support Table](CUDA_Runtime_API_functions_supported_by_HIP.md) for ### Is HIP a drop-in replacement for CUDA? No. HIP provides porting tools which do most of the work to convert CUDA code into portable C++ code that uses the HIP APIs. -Most developers will port their code from CUDA to HIP and then maintain the HIP version. +Most developers will port their code from CUDA to HIP and then maintain the HIP version. HIP code provides the same performance as native CUDA code, plus the benefits of running on AMD platforms. ### What specific version of CUDA does HIP support? -HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of the functionality provided in CUDA, and the hipify tools can -scan code to identify any unsupported CUDA functions - this is useful for identifying the specific features required by a given application. +HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of the functionality provided in CUDA, and the hipify tools can scan code to identify any unsupported CUDA functions - this is useful for identifying the specific features required by a given application. However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP. Each bullet below lists the major new language features in each CUDA release and then indicate which are supported/not supported in HIP: - CUDA 4.0 and earlier : - HIP supports CUDA 4.0 except for the limitations described above. - CUDA 5.0 : - - Dynamic Parallelism (not supported) + - Dynamic Parallelism (not supported) - cuIpc functions (under development). - CUDA 5.5 : - CUPTI (not directly supported, [AMD GPUPerfAPI](http://developer.amd.com/tools-and-sdks/graphics-development/gpuperfapi/) can be used as an alternative in some cases) @@ -98,10 +101,9 @@ However, we can provide a rough summary of the features included in each CUDA SD - CUDA 8.0 : - Page Migration including cudaMemAdvise, cudaMemPrefetch, other cudaMem* APIs(not supported) - ### What libraries does HIP support? HIP includes growing support for the four key math libraries using hcBlas, hcFft, hcrng and hcsparse, as well as MIOpen for machine intelligence applications. -These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaced with other HIP applications. +These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaced with other HIP applications. The hip interfaces support both ROCm and CUDA paths, with familiar library interfaces. - [hipBlas](https://github.com/ROCmSoftwarePlatform/hipBLAS), which utilizes [rocBlas](https://github.com/ROCmSoftwarePlatform/rocBLAS). @@ -110,7 +112,7 @@ The hip interfaces support both ROCm and CUDA paths, with familiar library inter - [hiprng](https://github.com/ROCmSoftwarePlatform/hcrng) Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the HIPIFY tools. These APIs use cublas or hcblas depending on the platform and replace the need -to use conditional compilation. +to use conditional compilation. ### How does HIP compare with OpenCL? Both AMD and Nvidia support OpenCL 1.2 on their devices so that developers can write portable code. @@ -131,7 +133,7 @@ HIP and CUDA provide similar math library calls as well. In summary, the HIP ph This reduces the potential for error, and also makes it easy to automate the translation. HIP's goal is to quickly get the ported program running on both platforms with little manual intervention, so that the programmer can focus on performance optimizations. -There have been several tools that have attempted to convert CUDA into OpenCL, such as CU2CL. OpenCL is a C99-based kernel language (rather than C++) and also does not support single-source compilation. +There have been several tools that have attempted to convert CUDA into OpenCL, such as CU2CL. OpenCL is a C99-based kernel language (rather than C++) and also does not support single-source compilation. As a result, the OpenCL syntax is different from CUDA, and the porting tools have to perform some heroic transformations to bridge this gap. The tools also struggle with more complex CUDA applications, in particular, those that use templates, classes, or other C++ features inside the kernel. @@ -153,52 +155,71 @@ NVCC is Nvidia's compiler driver for compiling "CUDA C++" code into PTX or devic ### What is HCC? HCC is AMD's compiler driver which compiles "heterogeneous C++" code into HSAIL or GCN device code for AMD GPUs. It's an open-source compiler based on recent versions of CLANG/LLVM. +In ROCM v3.5 release, HCC compiler is deprecated and HIP-Clang compiler is introduced to compile HIP programs. + +### What is HIP-Clang? +HIP-Clang is new compiler to emphasize its capability to compile HIP programs which can run on AMD platform. + ### Why use HIP rather than supporting CUDA directly? While HIP is a strong subset of the CUDA, it is a subset. The HIP layer allows that subset to be clearly defined and documented. -Developers who code to the HIP API can be assured their code will remain portable across Nvidia and AMD platforms. +Developers who code to the HIP API can be assured their code will remain portable across Nvidia and AMD platforms. In addition, HIP defines portable mechanisms to query architectural features and supports a larger 64-bit wavesize which expands the return type for cross-lane functions like ballot and shuffle from 32-bit ints to 64-bit ints. ### Can I develop HIP code on an Nvidia CUDA platform? -Yes. HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and HCC back-ends. -"Extra" APIs, parameters, and features which exist in CUDA but not in HCC will typically result in compile-time or run-time errors. +Yes. HIP's CUDA path only exposes the APIs and functionality that work on both NVCC and AMDGPU back-ends. +"Extra" APIs, parameters, and features which exist in CUDA but not in HIP-Clang will typically result in compile-time or run-time errors. Developers need to use the HIP API for most accelerator code and bracket any CUDA-specific code with preprocessor conditionals. Developers concerned about portability should, of course, run on both platforms, and should expect to tune for performance. In some cases, CUDA has a richer set of modes for some APIs, and some C++ capabilities such as virtual functions - see the HIP @API documentation for more details. -### Can I develop HIP code on an AMD HCC platform? -Yes. HIP's HCC path only exposes the APIs and functions that work on both NVCC and HCC back ends. "Extra" APIs, parameters and features that appear in HCC but not CUDA will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HCC-specific code with preprocessor conditionals. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HCC supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the hc path. +### Can I develop HIP code on an AMD HIP-Clang platform? +Yes. HIP's HIP-Clang path only exposes the APIs and functions that work on AMD runtime back ends. "Extra" APIs, parameters and features that appear in HIP-Clang but not CUDA will typically cause compile- or run-time errors. Developers must use the HIP API for most accelerator code and bracket any HIP-Clang specific code with preprocessor conditionals. Those concerned about portability should, of course, test their code on both platforms and should tune it for performance. Typically, HIP-Clang supports a more modern set of C++11/C++14/C++17 features, so HIP developers who want portability should be careful when using advanced C++ features on the HIP-Clang path. +In ROCM v3.5 release, HCC compiler is deprecated, and the HIP-Clang compiler can be used for compiling HIP programs. + +### Do I need to make code changes in HIP code if switching compiler from HCC to HIP-Clang? +For most HIP applications, the transition from HCC to HIP-Clang is transparent as the HIPCC and HIP cmake files automatically choose compiler options for HIP-Clang and hide the difference between the HCC and HIP-Clang code. +However, minor changes may be required as HIP-Clang has stricter syntax and semantic checks compared to HCC. + +### How to use HIP-Clang to build HIP programs? +The environment variable can be used to set compiler path: +- HIP_CLANG_PATH: path to hip-clang. When set, this variable let hipcc to use hip-clang for compilation/linking + +There is an alternative environment variable to set compiler path: +- HIP_ROCCLR_HOME: path to root directory of the HIP-ROCclr runtime. When set, this variable let hipcc use hip-clang from the ROCclr distribution. +NOTE: If HIP_ROCCLR_HOME is set, there is no need to set HIP_CLANG_PATH since hipcc will deduce them from HIP_ROCCLR_HOME. + +### What is ROCclr? +ROCclr (Radeon Open Compute Common Language Runtime) is a virtual device interface that compute runtimes interact with backends such as ROCr on Linux, as well as PAL on Windows. ### Can a HIP binary run on both AMD and Nvidia platforms? -HIP is a source-portable language that can be compiled to run on either the HCC or NVCC platform. HIP tools don't create a "fat binary" that can run on either platform, however. - +HIP is a source-portable language that can be compiled to run on either AMD or NVIDIA platform. HIP tools don't create a "fat binary" that can run on either platform, however. ### What's the difference between HIP and hc? HIP is a portable C++ language that supports a strong subset of the CUDA run-time APIs and device-kernel language. It's designed to simplify CUDA conversion to portable C++. HIP provides a C-compatible run-time API, C-compatible kernel-launch mechanism, C++ kernel language and pointer-based memory management. -A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard. +A C++ dialect, hc is supported by the AMD compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard. +The HCC compiler has been deprecated in the ROCm Release v3.5. - -### On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ? -Yes. HIP/HCC generates the object code which conforms to the GCC ABI, and also links with libstdc++. This means you can compile host code with the compiler of your choice and link the generated object code +### On HIP-Clang, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ? +Yes. HIP generates the object code which conforms to the GCC ABI, and also links with libstdc++. This means you can compile host code with the compiler of your choice and link the generated object code with GPU code compiled with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) and host code (compiled with gcc, icc, or clang). These projects can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from their preferred compiler. +### Can I install both CUDA SDK and HIP-Clang on the same machine? +Yes. You can use HIP_PLATFORM to choose which path hipcc targets. This configuration can be useful when using HIP to develop an application which is portable to both AMD and NVIDIA. -### HIP detected my platform (hcc vs nvcc) incorrectly - what should I do? -HIP will set the platform to HCC if it sees that the AMD graphics driver is installed and has detected an AMD GPU. -Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc) +### HIP detected my platform (HIP-Clang vs nvcc) incorrectly - what should I do? +HIP will set the platform to hcc and compiler to HIP-Clang if it sees that the AMD graphics driver is installed and has detected an AMD GPU. +Sometimes this isn't what you want - you can force HIP to recognize the platform by setting the following, ``` +export HIP_COMPILER=clang export HIP_PLATFORM=hcc ``` -One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56". This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc. HIP may be able to compile the application using the nvcc tool-chain but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild. - -If you see issues related to incorrect platform detection, please file an issue with the GitHub issue tracker so we can improve HIP's platform detection logic. - -### Can I install both CUDA SDK and HCC on the same machine? -Yes. You can use HIP_PLATFORM to choose which path hipcc targets. This configuration can be useful when using HIP to develop an application which is portable to both AMD and NVIDIA. +One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56". This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc. HIP may be able to compile the application using the nvcc tool-chain but will generate this error at runtime since the platform does not have a CUDA device. +The fix is to set HIP_PLATFORM=hcc and rebuild. ### On CUDA, can I mix CUDA code with HIP code? @@ -211,17 +232,8 @@ hipCUResultTohipError If platform portability is important, use #ifdef __HIP_PLATFORM_NVCC__ to guard the CUDA-specific code. -### On HCC, can I use HC functionality with HIP? -Yes. -The code can include hc.hpp and use HC functions inside the kernel. A typical use-case is to use AMD-specific hardware features such as the permute, swizzle, or DPP operations. -See the 'bit_extract' sample for an example. - -Also these functions can be used to extract HCC accelerator and accelerator_view structures from the HIP deviceId and hipStream_t: -hipHccGetAccelerator(int deviceId, hc::accelerator *acc); -hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av); - -If platform portability is important, use #ifdef __HIP_PLATFORM_HIPCC__ to guard the HCC-specific code. - +### On HIP-Clang, can I use HC functionality with HIP? +No. HC functionality is not supported by HIP-Clang. ### How do I trace HIP application flow? See the [HIP Profiling Guide](hip_porting_guide.md) for more information. @@ -231,12 +243,8 @@ Unlike CUDA, in HCC, for functions defined in the header files, the keyword of " Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation. The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__". -### How do I disable HIP Generic Grid Launch option? -Generic Grid Launch(GGL) is currently the default method for hip kernel launch. -To disable it and use the legancy grid launch method, please either change the default value of GENERIC_GRID_LAUNCH to 0 in the following to header files and rebuild HIP: -$HIP/include/hip/hcc_detail/hip_runtime_api.h -$HIP/include/hip/hcc_detail/host_defines.h -Or pass "-DGENERIC_GRID_LAUNCH=0" to hipcc at application compilation time. +### What is maximum limit of kernel launching parameter? +Product of block.x, block.y, and block.z should be less than 1024. -### What is maximum limit of Generic Grid Launch parameters (grid and block)? -Product of (grid.x and block.x), (grid.y and block.y) or (grid.z and block.z) should always be less than UINT_MAX. \ No newline at end of file +### Are __shfl_*_sync functions supported on HIP platform? +__shfl_*_sync is not supported on HIP but for nvcc path CUDA 9.0 and above all shuffle calls get redirected to it's sync version. diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index c291fa8ae6..816802fc79 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -1,13 +1,13 @@ # HIP Porting Guide In addition to providing a portable C++ programming environment for GPUs, HIP is designed to ease -the porting of existing CUDA code into the HIP environment. This section describes the available tools +the porting of existing CUDA code into the HIP environment. This section describes the available tools and provides practical suggestions on how to port CUDA code and work through common issues. ## Table of Contents -- [Porting a New Cuda Project](#porting-a-new-cuda-project) +- [Porting a New CUDA Project](#porting-a-new-cuda-project) * [General Tips](#general-tips) * [Scanning existing CUDA code to scope the porting effort](#scanning-existing-cuda-code-to-scope-the-porting-effort) * [Converting a project "in-place"](#converting-a-project-in-place) @@ -22,6 +22,7 @@ and provides practical suggestions on how to port CUDA code and work through com * [Device-Architecture Properties](#device-architecture-properties) * [Table of Architecture Properties](#table-of-architecture-properties) - [Finding HIP](#finding-hip) +- [Identifying HIP Runtime](#identifying-hip-runtime) - [hipLaunchKernel](#hiplaunchkernel) - [Compiler Options](#compiler-options) - [Linking Issues](#linking-issues) @@ -47,20 +48,20 @@ and provides practical suggestions on how to port CUDA code and work through com + [/usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr();](#usrincludecv1memory517215-error-call-to-implicitly-deleted-default-constructor-of-std__1bad_weak_ptr-throw-bad_weak_ptr) * [HIP Environment Variables](#hip-environment-variables) * [Editor Highlighting](#editor-highlighting) - + -## Porting a New Cuda Project +## Porting a New CUDA Project ### General Tips -- Starting the port on a Cuda machine is often the easiest approach, since you can incrementally port pieces of the code to HIP while leaving the rest in Cuda. (Recall that on Cuda machines HIP is just a thin layer over Cuda, so the two code types can interoperate on nvcc platforms.) Also, the HIP port can be compared with the original Cuda code for function and performance. -- Once the Cuda code is ported to HIP and is running on the Cuda machine, compile the HIP code using hcc on an AMD machine. -- HIP ports can replace Cuda versions: HIP can deliver the same performance as a native Cuda implementation, with the benefit of portability to both Nvidia and AMD architectures as well as a path to future C++ standard support. You can handle platform-specific features through conditional compilation or by adding them to the open-source HIP infrastructure. -- Use **[bin/hipconvertinplace-perl.sh](https://github.com/ROCm-Developer-Tools/HIP/blob/master/bin/hipconvertinplace-perl.sh)** to hipify all code files in the Cuda source directory. +- Starting the port on a CUDA machine is often the easiest approach, since you can incrementally port pieces of the code to HIP while leaving the rest in CUDA. (Recall that on CUDA machines HIP is just a thin layer over CUDA, so the two code types can interoperate on nvcc platforms.) Also, the HIP port can be compared with the original CUDA code for function and performance. +- Once the CUDA code is ported to HIP and is running on the CUDA machine, compile the HIP code using the HIP compiler on an AMD machine. +- HIP ports can replace CUDA versions: HIP can deliver the same performance as a native CUDA implementation, with the benefit of portability to both Nvidia and AMD architectures as well as a path to future C++ standard support. You can handle platform-specific features through conditional compilation or by adding them to the open-source HIP infrastructure. +- Use **[bin/hipconvertinplace-perl.sh](https://github.com/ROCm-Developer-Tools/HIP/blob/master/bin/hipconvertinplace-perl.sh)** to hipify all code files in the CUDA source directory. ### Scanning existing CUDA code to scope the porting effort -The hipexamine-perl.sh tool will scan a source directory to determine which files contain CUDA code and how much of that code can be automatically hipified, +The hipexamine-perl.sh tool will scan a source directory to determine which files contain CUDA code and how much of that code can be automatically hipified. ``` > cd examples/rodinia_3.0/cuda/kmeans > $HIP_DIR/bin/hipexamine-perl.sh. @@ -128,110 +129,119 @@ directory names. | CUDA Library | ROCm Library | Comment | |------- | --------- | ----- | -| cuBLAS | rocBLAS | Basic Linear Algebra Subroutines -| cuFFT | rocFFT | Fast Fourier Transfer Library -| cuSPARSE | rocSPARSE | Sparse BLAS + SPMV -| cuSolver | rocSolver | Lapack library +| cuBLAS | rocBLAS | Basic Linear Algebra Subroutines +| cuFFT | rocFFT | Fast Fourier Transfer Library +| cuSPARSE | rocSPARSE | Sparse BLAS + SPMV +| cuSolver | rocSOLVER | Lapack library | AMG-X | rocALUTION | Sparse iterative solvers and preconditioners with Geometric and Algebraic MultiGrid -| Thrust | hipThrust | C++ parallel algorithms library +| Thrust | rocThrust | C++ parallel algorithms library | CUB | rocPRIM | Low Level Optimized Parallel Primitives -| cuDNN | MIOpen | Deep learning Solver Library +| cuDNN | MIOpen | Deep learning Solver Library | cuRAND | rocRAND | Random Number Generator Library -| EIGEN | EIGEN – HIP port | C++ template library for linear algebra: matrices, vectors, numerical solvers, +| EIGEN | EIGEN – HIP port | C++ template library for linear algebra: matrices, vectors, numerical solvers, | NCCL | RCCL | Communications Primitives Library based on the MPI equivalents - + ## Distinguishing Compiler Modes - - + + ### Identifying HIP Target Platform -All HIP projects target either the hcc or nvcc platform. The platform affects which headers are included and which libraries are used for linking. - -- `HIPCC_PLATFORM_HCC` is defined if the HIP platform targets hcc -- `HIPCC_PLATFORM_NVCC` is defined if the HIP platform targets nvcc - -Many projects use a mixture of an accelerator compiler (hcc or nvcc) and a standard compiler (e.g., g++). These defines are set for both accelerator and standard compilers and thus are often the best option when writing code that uses conditional compilation. - - +All HIP projects target either AMD or NVIDIA platform. The platform affects which headers are included and which libraries are used for linking. + +- `HIP_PLATFORM_HCC` is defined if the HIP platform targets AMD + +- `HIP_PLATFORM_NVCC` is defined if the HIP platform targets NVIDIA + +On AMD platform, the compiler was hcc, but is deprecated in ROCM v3.5 release, and HIP-Clang compiler is introduced for compiling HIP programs. + +For most HIP applications, the transition from hcc to HIP-Clang is transparent. +HIPCC and HIP cmake files automatically choose compilation options for HIP-Clang and hide the difference between the hcc and hip-clang code. +However, minor changes may be required as HIP-Clang has stricter syntax and semantic checks compared to hcc. + +Many projects use a mixture of an accelerator compiler (AMD or NVIDIA) and a standard compiler (e.g. g++). These defines are set for both accelerator and standard compilers and thus are often the best option when writing code that uses conditional compilation. + + + ### Identifying the Compiler: hcc, hip-clang or nvcc -Often, it's useful to know whether the underlying compiler is hcc, hip-clang or nvcc. This knowledge can guard platform-specific code (features that only work on the nvcc, hip-clang or hcc path but not all) or aid in platform-specific performance tuning. - +Often, it's useful to know whether the underlying compiler is hcc, HIP-Clang or nvcc. This knowledge can guard platform-specific code or aid in platform-specific performance tuning. + + ``` #ifdef __HCC__ -// Compiled with hcc - +// Compiled with hcc + ``` ``` #ifdef __HIP__ -// Compiled with hip-clang - +// Compiled with HIP-Clang + ``` - + ``` #ifdef __NVCC__ -// Compiled with nvcc -// Could be compiling with Cuda language extensions enabled (for example, a ".cu file) +// Compiled with nvcc +// Could be compiling with CUDA language extensions enabled (for example, a ".cu file) // Could be in pass-through mode to an underlying host compile OR (for example, a .cpp file) - + ``` - + ``` #ifdef __CUDACC__ -// Compiled with nvcc (Cuda language extensions enabled) +// Compiled with nvcc (CUDA language extensions enabled) ``` - -hcc and hip-clang directly generates the host code (using the Clang x86 target) and passes the code to another host compiler. Thus, they have no equivalent of the \__CUDA_ACC define. - -The macro `__HIPCC__` is set if either `__HCC__`, `__HIP__` or `__CUDACC__` is defined. This configuration is useful in determining when code is being compiled using an accelerator-enabled compiler (hcc or nvcc) as opposed to a standard host compiler (GCC, ICC, Clang, etc.). - + +Compiler directly generates the host code (using the Clang x86 target) and passes the code to another host compiler. Thus, they have no equivalent of the \__CUDA_ACC define. + + ### Identifying Current Compilation Pass: Host or Device - -Both nvcc and hcc make two passes over the code: one for host code and one for device code. `__HIP_DEVICE_COMPILE__` is set to a nonzero value when the compiler (hcc or nvcc) is compiling code for a device inside a `__global__` kernel or for a device function. `__HIP_DEVICE_COMPILE__` can replace #ifdef checks on the `__CUDA_ARCH__` define. - + +nvcc makes two passes over the code: one for host code and one for device code. +HIP-Clang will have multiple passes over the code: one for the host code, and one for each architecture on the device code. +`__HIP_DEVICE_COMPILE__` is set to a nonzero value when the compiler (hcc, HIP-Clang or nvcc) is compiling code for a device inside a `__global__` kernel or for a device function. `__HIP_DEVICE_COMPILE__` can replace #ifdef checks on the `__CUDA_ARCH__` define. + ``` -// #ifdef __CUDA_ARCH__ +// #ifdef __CUDA_ARCH__ #if __HIP_DEVICE_COMPILE__ ``` - -Unlike `__CUDA_ARCH__`, the `__HIP_DEVICE_COMPILE__` value is 1 or undefined, and it doesn't represent the feature capability of the target device. + +Unlike `__CUDA_ARCH__`, the `__HIP_DEVICE_COMPILE__` value is 1 or undefined, and it doesn't represent the feature capability of the target device. ### Compiler Defines: Summary -|Define | hcc | hip-clang | nvcc | Other (GCC, ICC, Clang, etc.) +|Define | hcc | HIP-Clang | nvcc | Other (GCC, ICC, Clang, etc.) |--- | --- | --- | --- |---| |HIP-related defines:| |`__HIP_PLATFORM_HCC__`| Defined | Defined | Undefined | Defined if targeting hcc platform; undefined otherwise | |`__HIP_PLATFORM_NVCC__`| Undefined | Undefined | Defined | Defined if targeting nvcc platform; undefined otherwise | -|`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined +|`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined |`__HIPCC__` | Defined | Defined | Defined | Undefined -|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) |0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 +|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) |0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 |nvcc-related defines:| |`__CUDACC__` | Undefined | Undefined | Defined if source code is compiled by nvcc; undefined otherwise | Undefined |`__NVCC__` | Undefined | Undefined | Defined | Undefined -|`__CUDA_ARCH__` | Undefined | Undefined | Unsigned representing compute capability (e.g., "130") if in device code; 0 if in host code | Undefined +|`__CUDA_ARCH__` | Undefined | Undefined | Unsigned representing compute capability (e.g., "130") if in device code; 0 if in host code | Undefined |hcc-related defines:| |`__HCC__` | Defined | Undefined | Undefined | Undefined -|`__HCC_ACCELERATOR__` | Nonzero if in device code; otherwise undefined | Undefined | Undefined | Undefined +|`__HCC_ACCELERATOR__` | Nonzero if in device code; otherwise undefined | Undefined | Undefined | Undefined |hip-clang-related defines:| |`__HIP__` | Undefined | Defined | Undefined | Undefined -|hcc/hip-clang common defines:| +|hcc/HIP-Clang common defines:| |`__clang__` | Defined | Defined | Undefined | Defined if using Clang; otherwise undefined - ## Identifying Architecture Features ### HIP_ARCH Defines -Some Cuda code tests `__CUDA_ARCH__` for a specific value to determine whether the machine supports a certain architectural feature. For instance, +Some CUDA code tests `__CUDA_ARCH__` for a specific value to determine whether the machine supports a certain architectural feature. For instance, ``` -#if (__CUDA_ARCH__ >= 130) +#if (__CUDA_ARCH__ >= 130) // doubles are supported ``` -This type of code requires special attention, since hcc/AMD and nvcc/Cuda devices have different architectural capabilities. Moreover, you can't determine the presence of a feature using a simple comparison against an architecture's version number. HIP provides a set of defines and device properties to query whether a specific architectural feature is supported. +This type of code requires special attention, since hcc/AMD and nvcc/CUDA devices have different architectural capabilities. Moreover, you can't determine the presence of a feature using a simple comparison against an architecture's version number. HIP provides a set of defines and device properties to query whether a specific architectural feature is supported. -The `__HIP_ARCH_*` defines can replace comparisons of `__CUDA_ARCH__` values: +The `__HIP_ARCH_*` defines can replace comparisons of `__CUDA_ARCH__` values: ``` //#if (__CUDA_ARCH__ >= 130) // non-portable if __HIP_ARCH_HAS_DOUBLES__ { // portable HIP feature query @@ -281,7 +291,7 @@ The table below shows the full set of architectural properties that HIP supports |`__HIP_ARCH_HAS_SURFACE_FUNCS__` | hasSurfaceFuncs | |`__HIP_ARCH_HAS_3DGRID__` | has3dGrid | Grids and groups are 3D |`__HIP_ARCH_HAS_DYNAMIC_PARALLEL__` | hasDynamicParallelism | - + ## Finding HIP @@ -291,10 +301,24 @@ Makefiles can use the following syntax to conditionally provide a default HIP_PA HIP_PATH ?= $(shell hipconfig --path) ``` -## hipLaunchKernel +## Identifying HIP Runtime + +HIP can depend on ROCclr, or NVCC as runtime + +- AMD platform +`HIP_ROCclr` is defined on AMD platform that HIP use Radeon Open Compute Common Language Runtime, called ROCclr. + +ROCclr is a virtual device interface that HIP runtimes interact with different backends which allows runtimes to work on Linux , as well as Windows without much efforts. + +- NVIDIA platform +On Nvidia platform, HIP is just a thin layer on top of CUDA. +On non-AMD platform, HIP runtime determines if nvcc is available and can be used. If available, HIP_PLATFORM is set to nvcc and underneath CUDA path is used. + + +## hipLaunchKernel hipLaunchKernel is a variadic macro which accepts as parameters the launch configurations (grid dims, group dims, stream, dynamic shared size) followed by a variable number of kernel arguments. -This sequence is then expanded into the appropriate kernel launch syntax depending on the platform. +This sequence is then expanded into the appropriate kernel launch syntax depending on the platform. While this can be a convenient single-line kernel launch syntax, the macro implementation can cause issues when nested inside other macros. For example, consider the following: ``` @@ -325,57 +349,42 @@ MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCal ## Compiler Options -hipcc is a portable compiler driver that will call nvcc or hcc (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler. The `hipconfig` script may helpful in making -infrastructure that identifies the target platform and sets options appropriately. It returns either "nvcc" or "hcc." The following sample shows the script in a makefile: - -``` -HIP_PLATFORM=$(shell hipconfig --compiler) - -ifeq (${HIP_PLATFORM}, nvcc) - HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 -endif -ifeq (${HIP_PLATFORM}, hcc) - HIPCC_FLAGS = -Wno-deprecated-register -endif - -``` - +hipcc is a portable compiler driver that will call nvcc or HIP-Clang (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler. +The `hipconfig` script may helpful in identifying the target platform, compiler and runtime. It can also help set options appropriately. ## Linking Issues ### Linking With hipcc -hipcc adds the necessary libraries for HIP as well as for the accelerator compiler (nvcc or hcc). We recommend linking with hipcc. +hipcc adds the necessary libraries for HIP as well as for the accelerator compiler (nvcc or AMD compiler). We recommend linking with hipcc since it automatically links the binary to the necessary HIP runtime libraries. It also has knowledge on how to link and to manage the GPU objects. ### -lm Option - + hipcc adds -lm by default to the link command. ## Linking Code With Other Compilers -Cuda code often uses nvcc for accelerator code (defining and launching kernels, typically defined in .cu or .cuh files). -It also uses a standard compiler (g++) for the rest of the application. nvcc is a preprocessor that employs a standard host compiler (e.g., gcc) to generate the host code. -Code compiled using this tool can employ only the intersection of language features supported by both nvcc and the host compiler. -In some cases, you must take care to ensure the data types and alignment of the host compiler are identical to those of the device compiler. Only some host compilers are supported---for example, recent nvcc versions lack Clang host-compiler capability. +CUDA code often uses nvcc for accelerator code (defining and launching kernels, typically defined in .cu or .cuh files). +It also uses a standard compiler (g++) for the rest of the application. nvcc is a preprocessor that employs a standard host compiler (gcc) to generate the host code. +Code compiled using this tool can employ only the intersection of language features supported by both nvcc and the host compiler. +In some cases, you must take care to ensure the data types and alignment of the host compiler are identical to those of the device compiler. Only some host compilers are supported---for example, recent nvcc versions lack Clang host-compiler capability. hcc generates both device and host code using the same Clang-based compiler. The code uses the same API as gcc, which allows code generated by different gcc-compatible compilers to be linked together. For example, code compiled using hcc can link with code compiled using "standard" compilers (such as gcc, ICC and Clang). Take care to ensure all compilers use the same standard C++ header and library formats. ### libc++ and libstdc++ -Version 0.86 of hipcc now uses libstdc++ by default for the HCC platform. This improves cross-linking support between G++ and hcc, in particular for interfaces that use - standard C++ libraries (ie std::vector, std::string). +hipcc links to libstdc++ by default. This provides better compatibility between g++ and HIP. -If you pass "--stdlib=libc++" to hipcc, hipcc will use the libc++ library. Generally, libc++ provides a broader set of C++ features while libstdc++ is the standard -for more compilers (notably including g++). +If you pass "--stdlib=libc++" to hipcc, hipcc will use the libc++ library. Generally, libc++ provides a broader set of C++ features while libstdc++ is the standard for more compilers (notably including g++). -When cross-linking C++ code, any C++ functions that use types from the C++ standard library (including std::string, std::vector and other containers) must use the same standard-library implementation. They include the following: +When cross-linking C++ code, any C++ functions that use types from the C++ standard library (including std::string, std::vector and other containers) must use the same standard-library implementation. They include the following: - Functions or kernels defined in hcc that are called from a standard compiler - Functions defined in a standard compiler that are called from hcc. -Applications with these interfaces should use the default libstdc++ linking. +Applications with these interfaces should use the default libstdc++ linking. Applications which are compiled entirely with hipcc, and which benefit from advanced C++ features not supported in libstdc++, and which do not require portability to nvcc, may choose to use libc++. @@ -387,12 +396,12 @@ The hip_runtime.h and hip_runtime_api.h files define the types, functions and en - hip_runtime_api.h: defines all the HIP runtime APIs (e.g., hipMalloc) and the types required to call them. A source file that is only calling HIP APIs but neither defines nor launches any kernels can include hip_runtime_api.h. hip_runtime_api.h uses no custom hc language features and can be compiled using a standard C++ compiler. - hip_runtime.h: included in hip_runtime_api.h. It additionally provides the types and defines required to create and launch kernels. hip_runtime.h does use custom hc language features, but they are guarded by ifdef checks. It can be compiled using a standard C++ compiler but will expose a subset of the available functions. -Cuda has slightly different contents for these two files. In some cases you may need to convert hipified code to include the richer hip_runtime.h instead of hip_runtime_api.h. +CUDA has slightly different contents for these two files. In some cases you may need to convert hipified code to include the richer hip_runtime.h instead of hip_runtime_api.h. ### Using a Standard C++ Compiler You can compile hip\_runtime\_api.h using a standard C or C++ compiler (e.g., gcc or ICC). The HIP include paths and defines (`__HIP_PLATFORM_HCC__` or `__HIP_PLATFORM_NVCC__`) must pass to the standard compiler; hipconfig then returns the necessary options: ``` -> hipconfig --cxx_config +> hipconfig --cxx_config -D__HIP_PLATFORM_HCC__ -I/home/user1/hip/include ``` @@ -410,11 +419,11 @@ The hipify-perl script automatically converts "cuda_runtime.h" to "hip_runtime.h #### cuda.h -The hcc path provides an empty cuda.h file. Some existing Cuda programs include this file but don't require any of the functions. +The hcc path provides an empty cuda.h file. Some existing CUDA programs include this file but don't require any of the functions. ### Choosing HIP File Extensions -Many existing Cuda projects use the ".cu" and ".cuh" file extensions to indicate code that should be run through the nvcc compiler. +Many existing CUDA projects use the ".cu" and ".cuh" file extensions to indicate code that should be run through the nvcc compiler. For quick HIP ports, leaving these file extensions unchanged is often easier, as it minimizes the work required to change file names in the directory and #include statements in the files. For new projects or ports which can be re-factored, we recommend the use of the extension ".hip.cpp" for source files, and @@ -547,7 +556,7 @@ hipcc-cmd: /opt/hcc/bin/hcc -hc -I/opt/hcc/include -stdlib=libc++ -I../../../.. #### /usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr(); -If you pass a ".cu" file, hcc will attempt to compile it as a Cuda language file. You must tell hcc that it's in fact a C++ file: use the "-x c++" option. +If you pass a ".cu" file, hcc will attempt to compile it as a CUDA language file. You must tell hcc that it's in fact a C++ file: use the "-x c++" option. ### HIP Environment Variables diff --git a/docs/markdown/hip_terms2.md b/docs/markdown/hip_terms2.md index 3b4661729d..8065c47876 100644 --- a/docs/markdown/hip_terms2.md +++ b/docs/markdown/hip_terms2.md @@ -10,10 +10,16 @@ The default device can be set with hipSetDevice. - completion_future becomes ready. "Completes". -- hcc = Heterogeneous Compute Compiler (https://bitbucket.org/multicoreware/hcc/wiki/Home). +- hcc = Heterogeneous Compute Compiler ( https://github.com/RadeonOpenCompute/hcc). + Starting from ROCM v3.5 release, hcc compiler is deprecated and HIP-Clang compiler is introduced for compiling HIP programs + +- HIP-Clang - Heterogeneous AMDGPU Compiler, with its capability to compile HIP programs on AMD platform (https://github.com/RadeonOpenCompute/llvm-project). + +- ROCclr - a virtual device interface that compute runtimes interact with different backends such as ROCr on Linux or PAL on Windows. + The ROCclr (https://github.com/ROCm-Developer-Tools/ROCclr) is an abstraction layer allowing runtimes to work on both OSes without much effort. - hipify tools - tools to convert CUDA(R) code to portable C++ code (https://github.com/ROCm-Developer-Tools/HIPIFY). + - hipconfig - tool to report various configuration properties of the target platform. - nvcc = nvcc compiler, do not capitalize. -- hcc = heterogeneous compute compiler, do not capitalize. diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d908f72a04..df5c4e5b60 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1375,10 +1375,12 @@ hipError_t hipHostFree(void* ptr); * * For hipMemcpy, the copy is always performed by the current device (set by hipSetDevice). * For multi-gpu or peer-to-peer configurations, it is recommended to set the current device to the - * device where the src data is physically located. For optimal peer-to-peer copies, the copy device - * must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy - * agent as the current device and src/dest as the peerDevice argument. if this is not done, the - * hipMemcpy will still work, but will perform the copy using a staging buffer on the host. + * device where the src data is physically located. For optimal peer-to-peer copies, the copy device + * must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy + * agent as the current device and src/dest as the peerDevice argument. if this is not done, the + * hipMemcpy will still work, but will perform the copy using a staging buffer on the host. + * Calling hipMemcpy with dst and src pointers that do not match the hipMemcpyKind results in + * undefined behavior. * * @param[out] dst Data being copy to * @param[in] src Data being copy from @@ -2089,13 +2091,14 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp /** * @brief Copies data between host and device. * - * @param[in] dst Destination memory address - * @param[in] dpitch Pitch of destination memory - * @param[in] src Source memory address - * @param[in] spitch Pitch of source memory - * @param[in] width Width of matrix transfer (columns in bytes) - * @param[in] height Height of matrix transfer (rows) - * @param[in] kind Type of transfer + * @param[in] dst Destination memory address + * @param[in] wOffset Destination starting X offset + * @param[in] hOffset Destination starting Y offset + * @param[in] src Source memory address + * @param[in] spitch Pitch of source memory + * @param[in] width Width of matrix transfer (columns in bytes) + * @param[in] height Height of matrix transfer (rows) + * @param[in] kind Type of transfer * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection * @@ -2108,13 +2111,12 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con /** * @brief Copies data between host and device. * - * @param[in] dst Destination memory address - * @param[in] dpitch Pitch of destination memory - * @param[in] src Source memory address - * @param[in] spitch Pitch of source memory - * @param[in] width Width of matrix transfer (columns in bytes) - * @param[in] height Height of matrix transfer (rows) - * @param[in] kind Type of transfer + * @param[in] dst Destination memory address + * @param[in] wOffset Destination starting X offset + * @param[in] hOffset Destination starting Y offset + * @param[in] src Source memory address + * @param[in] count size in bytes to copy + * @param[in] kind Type of transfer * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection * @@ -3375,7 +3377,7 @@ hipError_t hipLaunchKernel(const void* function_address, size_t sharedMemBytes __dparm(0), hipStream_t stream __dparm(0)); -#if __HIP_ROCclr__ +#if __HIP_ROCclr__ || !defined(__HCC__) hipError_t hipBindTexture( size_t* offset, const textureReference* tex, @@ -3654,6 +3656,7 @@ hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg); hipError_t hipRemoveActivityCallback(uint32_t id); const char* hipApiName(uint32_t id); const char* hipKernelNameRef(const hipFunction_t f); +const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream); #ifdef __cplusplus } /* extern "C" */ #endif diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 23b8a0619d..fe72f33d65 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -236,6 +236,13 @@ typedef struct cudaResourceViewDesc hipResourceViewDesc; #define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT #define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX +#if CUDA_VERSION >= 9000 +#define __shfl(...) __shfl_sync(0xffffffff, __VA_ARGS__) +#define __shfl_up(...) __shfl_up_sync(0xffffffff, __VA_ARGS__) +#define __shfl_down(...) __shfl_down_sync(0xffffffff, __VA_ARGS__) +#define __shfl_xor(...) __shfl_xor_sync(0xffffffff, __VA_ARGS__) +#endif // CUDA_VERSION >= 9000 + inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { switch (cuError) { case cudaSuccess: diff --git a/packaging/hip-base.txt b/packaging/hip-base.txt index 0923f0c8fd..6a5147c292 100644 --- a/packaging/hip-base.txt +++ b/packaging/hip-base.txt @@ -29,7 +29,7 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR} set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "perl (>= 5.0)") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "perl (>= 5.0),libfile-which-perl") set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip-base") set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_base") set(CPACK_BINARY_RPM "ON") @@ -37,7 +37,7 @@ set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") -set(CPACK_RPM_PACKAGE_REQUIRES "perl >= 5.0") +set(CPACK_RPM_PACKAGE_REQUIRES "perl >= 5.0,perl-File-Which") set(CPACK_RPM_PACKAGE_OBSOLETES "hip_base") set(CPACK_RPM_PACKAGE_CONFLICTS "hip_base") set(CPACK_BINARY_RPM "ON") diff --git a/packaging/hip-rocclr.txt b/packaging/hip-rocclr.txt index 04eb892cc8..c9221a64c6 100644 --- a/packaging/hip-rocclr.txt +++ b/packaging/hip-rocclr.txt @@ -40,7 +40,7 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR} set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "hsa-rocr-dev, hsa-ext-rocr-dev, rocm-utils, hip-base (= ${CPACK_PACKAGE_VERSION}), comgr (>= 1.1), llvm-amdgpu") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hsa-rocr-dev, rocm-utils, hip-base (= ${CPACK_PACKAGE_VERSION}), comgr (>= 1.1), llvm-amdgpu") set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_rocclr, hip-hcc (= ${CPACK_PACKAGE_VERSION})") set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_rocclr") set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_rocclr") @@ -50,7 +50,7 @@ set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) -set(CPACK_RPM_PACKAGE_REQUIRES "hsa-rocr-dev, hsa-ext-rocr-dev, rocm-utils, hip-base = ${HIP_BASE_VERSION}, comgr >= 1.1, llvm-amdgpu") +set(CPACK_RPM_PACKAGE_REQUIRES "hsa-rocr-dev, rocm-utils, hip-base = ${HIP_BASE_VERSION}, comgr >= 1.1, llvm-amdgpu") set(CPACK_RPM_PACKAGE_PROVIDES "hip_rocclr, hip-hcc = ${HIP_BASE_VERSION}") set(CPACK_RPM_PACKAGE_OBSOLETES "hip_rocclr") set(CPACK_RPM_PACKAGE_CONFLICTS "hip_rocclr") diff --git a/rocclr/hip_device_runtime.cpp b/rocclr/hip_device_runtime.cpp index 3823af5206..9ed1c64145 100644 --- a/rocclr/hip_device_runtime.cpp +++ b/rocclr/hip_device_runtime.cpp @@ -442,6 +442,9 @@ hipError_t hipDeviceSynchronize ( void ) { } queue->finish(); + + hip::Stream::syncNonBlockingStreams(); + HIP_RETURN(hipSuccess); } diff --git a/rocclr/hip_event.cpp b/rocclr/hip_event.cpp index 2e4870834f..584a133613 100644 --- a/rocclr/hip_event.cpp +++ b/rocclr/hip_event.cpp @@ -90,11 +90,22 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) { return hipErrorNotReady; } - // For certain HIP Api's that take start and stop event - // the command is the same - if (event_ == eStop.event_) { + // For certain HIP API's that take start and stop event + // and no hipEventRecord needs to be called + if (event_ == eStop.event_ && !recorded_ && !eStop.recorded_) { ms = static_cast(static_cast(eStop.event_->profilingInfo().end_ - event_->profilingInfo().start_))/1000000.f; + } else if (event_ == eStop.event_) { + // Events are the same, which indicates the stream is empty and likely + // eventRecord is called on another stream. For such cases insert and measure a + // marker. + amd::Command* command = new amd::Marker(*event_->command().queue(), false); + command->enqueue(); + command->awaitCompletion(); + ms = static_cast(static_cast(command->event().profilingInfo().end_ - + event_->profilingInfo().end_))/1000000.f; + command->release(); + } else { ms = static_cast(static_cast(eStop.event_->profilingInfo().end_ - event_->profilingInfo().end_))/1000000.f; @@ -126,7 +137,7 @@ hipError_t Event::streamWait(amd::HostQueue* hostQueue, uint flags) { return hipSuccess; } -void Event::addMarker(amd::HostQueue* queue, amd::Command* command) { +void Event::addMarker(amd::HostQueue* queue, amd::Command* command, bool record) { amd::ScopedLock lock(lock_); if (event_ == &command->event()) return; @@ -136,6 +147,7 @@ void Event::addMarker(amd::HostQueue* queue, amd::Command* command) { } event_ = &command->event(); + recorded_ = record; } } @@ -233,7 +245,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { } hip::Event* e = reinterpret_cast(event); - e->addMarker(queue, command); + e->addMarker(queue, command, true); HIP_RETURN(hipSuccess); } diff --git a/rocclr/hip_event.hpp b/rocclr/hip_event.hpp index 2360c972bb..78e081d4dc 100644 --- a/rocclr/hip_event.hpp +++ b/rocclr/hip_event.hpp @@ -37,7 +37,8 @@ public: class Event { public: - Event(unsigned int flags) : flags(flags), lock_("hipEvent_t"), event_(nullptr) { + Event(unsigned int flags) : flags(flags), lock_("hipEvent_t"), + event_(nullptr), recorded_(false) { // No need to init event_ here as addMarker does that } @@ -53,13 +54,18 @@ public: hipError_t elapsedTime(Event& stop, float& ms); hipError_t streamWait(amd::HostQueue* queue, uint flags); - void addMarker(amd::HostQueue* queue, amd::Command* command); + void addMarker(amd::HostQueue* queue, amd::Command* command, bool record); private: amd::Monitor lock_; amd::HostQueue* stream_; amd::Event* event_; + //! Flag to indicate hipEventRecord has been called. This is needed except for + //! hip*ModuleLaunchKernel API which takes start and stop events so no + //! hipEventRecord is called. Cleanup needed once those APIs are deprecated. + bool recorded_; + bool ready(); }; diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in index 53bdfb3065..a3476c7f33 100755 --- a/rocclr/hip_hcc.def.in +++ b/rocclr/hip_hcc.def.in @@ -253,3 +253,4 @@ hipTexObjectGetResourceDesc hipTexObjectGetResourceViewDesc hipTexObjectGetTextureDesc hipExtStreamCreateWithCUMask +hipStreamGetPriority diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 02a704bf88..11637f2696 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -181,6 +181,7 @@ global: hipRemoveActivityCallback; hipApiName; hipKernelNameRef; + hipKernelNameRefByPtr; hipProfilerStart; hipProfilerStop; hiprtcCompileProgram; @@ -259,6 +260,7 @@ global: hipEnableActivityCallback*; hipGetCmdName*; hipExtStreamCreateWithCUMask; + hipStreamGetPriority; }; local: *; diff --git a/rocclr/hip_intercept.cpp b/rocclr/hip_intercept.cpp index f2e58b856b..aa61d6bc38 100644 --- a/rocclr/hip_intercept.cpp +++ b/rocclr/hip_intercept.cpp @@ -19,6 +19,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" +#include "hip_internal.hpp" #include "hip_prof_api.h" // HIP API callback/activity @@ -27,6 +28,19 @@ api_callbacks_table_t callbacks_table; extern const std::string& FunctionName(const hipFunction_t f); const char* hipKernelNameRef(const hipFunction_t f) { return FunctionName(f).c_str(); } +const char* hipKernelNameRefByPtr(const void *hostFunction, hipStream_t stream) { + hip::Stream* s = reinterpret_cast(stream); + int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); + if (deviceId == -1) { + DevLogPrintfError("Wrong Device Id: %d \n", deviceId); + return NULL; + } + hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); + if (func == nullptr) { + return NULL; + } + return hipKernelNameRef(func); +} hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) { return callbacks_table.set_callback(id, reinterpret_cast(fun), arg) ? diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 492842c32b..702b93c8b3 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -109,6 +109,11 @@ namespace hip { amd::Monitor& Lock() const { return lock_; } /// Returns the creation flags for the current stream unsigned int Flags() const { return flags_; } + /// Returns the priority for the current stream + amd::CommandQueue::Priority Priority() const { return priority_; } + + /// Sync all non-blocking streams + static void syncNonBlockingStreams(); }; /// HIP Device class diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index b7b8d77592..290b98cd45 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -761,7 +761,7 @@ hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeByt if ((offset + sizeBytes) > sym_size) { DevLogPrintfError("Trying to access out of bounds, offset: %u sizeBytes: %u sym_size: %u \n", offset, sizeBytes, sym_size); - return HIP_RETURN(hipErrorInvalidDevicePointer); + HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -794,7 +794,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t sizeBytes, if ((offset + sizeBytes) > sym_size) { DevLogPrintfError("Trying to access out of bounds, offset: %u sizeBytes: %u sym_size: %u \n", offset, sizeBytes, sym_size); - return HIP_RETURN(hipErrorInvalidDevicePointer); + HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -827,7 +827,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si if ((offset + sizeBytes) > sym_size) { DevLogPrintfError("Trying to access out of bounds, offset: %u sizeBytes: %u sym_size: %u \n", offset, sizeBytes, sym_size); - return HIP_RETURN(hipErrorInvalidDevicePointer); + HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -860,7 +860,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBy if ((offset + sizeBytes) > sym_size) { DevLogPrintfError("Trying to access out of bounds, offset: %u sizeBytes: %u sym_size: %u \n", offset, sizeBytes, sym_size); - return HIP_RETURN(hipErrorInvalidDevicePointer); + HIP_RETURN(hipErrorInvalidDevicePointer); } device_ptr = reinterpret_cast
(device_ptr) + offset; @@ -1737,28 +1737,31 @@ hipError_t ihipMemset(void* dst, int64_t value, size_t valueSize, size_t sizeByt int64_t value64 = 0; if (sizeBytes/sizeof(int64_t) > 0) { n_head_bytes = static_cast(aligned_dst) - static_cast(dst); - if (valueSize == sizeof(int8_t)) { - value = value & 0xff; - value64 = ((value << 56) | (value << 48) | (value << 40) | (value << 32) - | (value << 24) | (value << 16) | (value << 8) | (value)); - } else if (valueSize == sizeof(int16_t)) { - value = value & 0xffff; - value64 = ((value << 48) | (value << 32) | (value << 16) | (value)); - } else if (valueSize == sizeof(int32_t)) { - value = value & 0xffffffff; - value64 = ((value << 32) | (value)); - } else if (valueSize == sizeof(int64_t)) { - value64 = value; - } else { - LogPrintfError("Unsupported Pattern size: %u \n", valueSize); - return hipErrorInvalidValue; - } n_tail_bytes = ((sizeBytes - n_head_bytes) % sizeof(int64_t)); - // If n_tail_bytes is != 0 then we will do a second fillBuffer Command - // on the same stream below, dont wait, do the first call async. - hip_error = packFillMemoryCommand(memory, offset, value64, sizeof(int64_t), - sizeBytes - n_tail_bytes - n_head_bytes, queue, - ((n_head_bytes != 0) || (n_tail_bytes != 0) || isAsync)); + size_t n_bytes = sizeBytes - n_tail_bytes - n_head_bytes; + if (n_bytes > 0) { + if (valueSize == sizeof(int8_t)) { + value = value & 0xff; + value64 = ((value << 56) | (value << 48) | (value << 40) | (value << 32) + | (value << 24) | (value << 16) | (value << 8) | (value)); + } else if (valueSize == sizeof(int16_t)) { + value = value & 0xffff; + value64 = ((value << 48) | (value << 32) | (value << 16) | (value)); + } else if (valueSize == sizeof(int32_t)) { + value = value & 0xffffffff; + value64 = ((value << 32) | (value)); + } else if (valueSize == sizeof(int64_t)) { + value64 = value; + } else { + LogPrintfError("Unsupported Pattern size: %u \n", valueSize); + return hipErrorInvalidValue; + } + // If n_tail_bytes is != 0 then we will do a second fillBuffer Command + // on the same stream below, dont wait, do the first call async. + hip_error = packFillMemoryCommand(memory, offset, value64, sizeof(int64_t), + n_bytes, queue, + ((n_head_bytes != 0) || (n_tail_bytes != 0) || isAsync)); + } if (hip_error != hipSuccess) { return hip_error; } diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 3a51001063..8f3d4ca936 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -392,11 +392,12 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, return hipErrorLaunchFailure; } int num_blocks = 0; - int num_grids = 0; + int max_blocks_per_grid = 0; + int best_block_size = 0; int block_size = blockDimX * blockDimY * blockDimZ; hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, &num_grids, device, f, block_size, sharedMemBytes, true); - if (((gridDimX * gridDimY * gridDimZ) / block_size) > unsigned(num_grids)) { + &num_blocks, &max_blocks_per_grid, &best_block_size, device, f, block_size, sharedMemBytes, true); + if (((gridDimX * gridDimY * gridDimZ) / block_size) > unsigned(max_blocks_per_grid)) { return hipErrorCooperativeLaunchTooLarge; } } @@ -458,11 +459,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, command->enqueue(); if(startEvent != nullptr) { - eStart->addMarker(queue, command); + eStart->addMarker(queue, command, false); command->retain(); } if(stopEvent != nullptr) { - eStop->addMarker(queue, command); + eStop->addMarker(queue, command, false); command->retain(); } command->release(); diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 0be928f9a4..d23197bd9e 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -125,7 +125,7 @@ hipError_t __hipExtractCodeObjectFromFatBinary(const void* data, if (num_code_objs == devices.size()) { return hipSuccess; } else { - DevLogError("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); + fatal("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); return hipErrorNoBinaryForGpu; } } @@ -442,10 +442,10 @@ hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) { if (createFunc(&function, module, devFunc.deviceName.c_str()) && function != nullptr) { devFunc.functions[deviceId] = function; - } - else { - // tprintf(DB_FB, "__hipRegisterFunction cannot find kernel %s for" - // " device %d\n", deviceName, deviceId); + } else { + DevLogPrintfError("__hipRegisterFunction cannot find kernel %s for device %d\n", + devFunc.deviceName.c_str(), deviceId); + return nullptr; } } return devFunc.functions[deviceId]; @@ -547,7 +547,9 @@ bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t dvar->rvars[deviceId].amd_mem_obj_ = amd_mem_obj; amd::MemObjMap::AddMemObj(device_ptr, amd_mem_obj); } else { - LogError("__hipRegisterVar cannot find kernel for device \n"); + DevLogPrintfError("__hipRegisterVar cannot find Var: %s for deviceId: 0x%x \n", + dvar->hostVar.c_str(), deviceId); + return false; } } *size_ptr = dvar->rvars[deviceId].getvarsize(); @@ -855,31 +857,37 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor namespace hip_impl { hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, int* numGrids, - const amd::Device& device, hipFunction_t func, int blockSize, + int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize, + const amd::Device& device, hipFunction_t func, int inputBlockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz) { hip::Function* function = hip::Function::asFunction(func); const amd::Kernel& kernel = *function->function_; const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel.getDeviceKernel(device)->workGroupInfo(); - if (blockSize == 0) { - if (bCalcPotentialBlkSz == false){ + if (bCalcPotentialBlkSz == false) { + if (inputBlockSize == 0) { return hipErrorInvalidValue; } - else { - blockSize = device.info().maxWorkGroupSize_; // maxwavefrontperblock + *bestBlockSize = 0; + // Make sure the requested block size is smaller than max supported + if (inputBlockSize > int(device.info().maxWorkGroupSize_)) { + *maxBlocksPerCU = 0; + *numBlocksPerGrid = 0; + return hipSuccess; } } - - // Make sure the requested block size is smaller than max supported - if (blockSize > int(device.info().maxWorkGroupSize_)) { - numBlocks = 0; - numGrids = 0; - return hipSuccess; + else { + if (inputBlockSize > int(device.info().maxWorkGroupSize_) || + inputBlockSize == 0) { + // The user wrote the kernel to work with a workgroup size + // bigger than this hardware can support. Or they do not care + // about the size So just assume its maximum size is + // constrained by hardware + inputBlockSize = device.info().maxWorkGroupSize_; + } } - - // Find threads accupancy per CU => simd_per_cu * GPR usage + // Find wave occupancy per CU => simd_per_cu * GPR usage constexpr size_t MaxWavesPerSimd = 8; // Limited by SPI 32 per CU, hence 8 per SIMD size_t VgprWaves = MaxWavesPerSimd; if (wrkGrpInfo->usedVGPRs_ > 0) { @@ -888,26 +896,49 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( size_t GprWaves = VgprWaves; if (wrkGrpInfo->usedSGPRs_ > 0) { - const size_t maxSGPRs = (device.info().gfxipVersion_ < 800) ? 512 : 800; - size_t SgprWaves = maxSGPRs / amd::alignUp(wrkGrpInfo->usedSGPRs_, 16); + size_t maxSGPRs; + if (device.info().gfxipVersion_ < 800) { + maxSGPRs = 512; + } + else if (device.info().gfxipVersion_ < 1000) { + maxSGPRs = 800; + } + else { + maxSGPRs = SIZE_MAX; // gfx10+ does not share SGPRs between waves + } + const size_t SgprWaves = maxSGPRs / amd::alignUp(wrkGrpInfo->usedSGPRs_, 16); GprWaves = std::min(VgprWaves, SgprWaves); } - size_t alu_accupancy = device.info().simdPerCU_ * std::min(MaxWavesPerSimd, GprWaves); - alu_accupancy *= wrkGrpInfo->wavefrontSize_; - // Calculate blocks occupancy per CU - *numBlocks = alu_accupancy / amd::alignUp(blockSize, wrkGrpInfo->wavefrontSize_); + const size_t alu_occupancy = device.info().simdPerCU_ * std::min(MaxWavesPerSimd, GprWaves); + const int alu_limited_threads = alu_occupancy * wrkGrpInfo->wavefrontSize_; - size_t total_used_lds = wrkGrpInfo->usedLDSSize_ + dynamicSMemSize; + int lds_occupancy_wgs = INT_MAX; + const size_t total_used_lds = wrkGrpInfo->usedLDSSize_ + dynamicSMemSize; if (total_used_lds != 0) { - // Calculate LDS occupancy per CU. lds_per_cu / (static_lsd + dynamic_lds) - int lds_occupancy = static_cast(device.info().localMemSize_ / total_used_lds); - *numBlocks = std::min(*numBlocks, lds_occupancy); + lds_occupancy_wgs = static_cast(device.info().localMemSize_ / total_used_lds); } + // Calculate how many blocks of inputBlockSize we can fit per CU + // Need to align with hardware wavefront size. If they want 65 threads, but + // waves are 64, then we need 128 threads per block. + // So this calculates how many blocks we can fit. + *maxBlocksPerCU = alu_limited_threads / amd::alignUp(inputBlockSize, wrkGrpInfo->wavefrontSize_); + // Unless those blocks are further constrained by LDS size. + *maxBlocksPerCU = std::min(*maxBlocksPerCU, lds_occupancy_wgs); - if (bCalcPotentialBlkSz) { - *numGrids = *numBlocks * device.info().numRTCUs_; - } + // Some callers of this function want to return the block size, in threads, that + // leads to the maximum occupancy. In that case, inputBlockSize is the maximum + // workgroup size the user wants to allow, or that the hardware can allow. + // It is either the number of threads that we are limited to due to occupancy, or + // the maximum available block size for this kernel, which could have come from the + // user. e.g., if the user indicates the maximum block size is 64 threads, but we + // calculate that 128 threads can fit in each CU, we have to give up and return 64. + *bestBlockSize = std::min(alu_limited_threads, amd::alignUp(inputBlockSize, wrkGrpInfo->wavefrontSize_)); + // If the best block size is smaller than the block size used to fit the maximum, + // then we need to make the grid bigger for full occupancy. + const int bestBlocksPerCU = alu_limited_threads / (*bestBlockSize); + // Unless those blocks are further constrained by LDS size. + *numBlocksPerGrid = device.info().maxComputeUnits_ * std::min(bestBlocksPerCU, lds_occupancy_wgs); return hipSuccess; } @@ -920,20 +951,21 @@ hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, { HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, f, dynSharedMemPerBlk, blockSizeLimit); if ((gridSize == nullptr) || (blockSize == nullptr)) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); if (func == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; - int num_grids = 0; + int max_blocks_per_grid = 0; int num_blocks = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, &num_grids, device, func, 0, dynSharedMemPerBlk,true); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, func, blockSizeLimit, dynSharedMemPerBlk,true); if (ret == hipSuccess) { - *blockSize = num_blocks; - *gridSize = num_grids; + *blockSize = best_block_size; + *gridSize = max_blocks_per_grid; } HIP_RETURN(ret); } @@ -944,16 +976,17 @@ hipError_t hipModuleOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize { HIP_INIT_API(hipModuleOccupancyMaxPotentialBlockSize, f, dynSharedMemPerBlk, blockSizeLimit); if ((gridSize == nullptr) || (blockSize == nullptr)) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; - int num_grids = 0; + int max_blocks_per_grid = 0; int num_blocks = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, &num_grids, device, f, 0, dynSharedMemPerBlk,true); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, f, blockSizeLimit, dynSharedMemPerBlk,true); if (ret == hipSuccess) { - *blockSize = num_blocks; - *gridSize = num_grids; + *blockSize = best_block_size; + *gridSize = max_blocks_per_grid; } HIP_RETURN(ret); } @@ -964,16 +997,17 @@ hipError_t hipModuleOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int* { HIP_INIT_API(hipModuleOccupancyMaxPotentialBlockSizeWithFlags, f, dynSharedMemPerBlk, blockSizeLimit, flags); if ((gridSize == nullptr) || (blockSize == nullptr)) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; - int num_grids = 0; + int max_blocks_per_grid = 0; int num_blocks = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, &num_grids, device, f, 0, dynSharedMemPerBlk,true); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, f, blockSizeLimit, dynSharedMemPerBlk,true); if (ret == hipSuccess) { - *blockSize = num_blocks; - *gridSize = num_grids; + *blockSize = best_block_size; + *gridSize = max_blocks_per_grid; } HIP_RETURN(ret); } @@ -983,13 +1017,15 @@ hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, { HIP_INIT_API(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor, f, blockSize, dynSharedMemPerBlk); if (numBlocks == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; int num_blocks = 0; + int max_blocks_per_grid = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, nullptr, device, f, blockSize, dynSharedMemPerBlk, false); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, f, blockSize, dynSharedMemPerBlk, false); *numBlocks = num_blocks; HIP_RETURN(ret); } @@ -1000,13 +1036,15 @@ hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numB { HIP_INIT_API(hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, f, blockSize, dynSharedMemPerBlk, flags); if (numBlocks == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; int num_blocks = 0; + int max_blocks_per_grid = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, nullptr, device, f, blockSize, dynSharedMemPerBlk, false); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, f, blockSize, dynSharedMemPerBlk, false); *numBlocks = num_blocks; HIP_RETURN(ret); } @@ -1016,19 +1054,21 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, { HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, f, blockSize, dynamicSMemSize); if (numBlocks == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); if (func == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; int num_blocks = 0; + int max_blocks_per_grid = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, nullptr, device, func, blockSize, dynamicSMemSize, false); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, func, blockSize, dynamicSMemSize, false); *numBlocks = num_blocks; HIP_RETURN(ret); } @@ -1039,19 +1079,21 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, { HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, f, blockSize, dynamicSMemSize, flags); if (numBlocks == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); if (func == nullptr) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; int num_blocks = 0; + int max_blocks_per_grid = 0; + int best_block_size = 0; hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( - &num_blocks, nullptr, device, func, blockSize, dynamicSMemSize, false); + &num_blocks, &max_blocks_per_grid, &best_block_size, device, func, blockSize, dynamicSMemSize, false); *numBlocks = num_blocks; HIP_RETURN(ret); } diff --git a/rocclr/hip_platform.hpp b/rocclr/hip_platform.hpp index 8e5eaa191f..fcbfb53bbb 100644 --- a/rocclr/hip_platform.hpp +++ b/rocclr/hip_platform.hpp @@ -23,7 +23,7 @@ namespace hip_impl { hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, int* numGrids, + int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize, const amd::Device& device, hipFunction_t func, int blockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz); -} \ No newline at end of file +} diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp old mode 100644 new mode 100755 index 23b89e71ca..43fce39299 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -57,10 +57,8 @@ bool Stream::Create() { bool result = (queue_ != nullptr) ? queue_->create() : false; // Insert just created stream into the list of the blocking queues if (result) { - if (!(flags_ & hipStreamNonBlocking)) { - amd::ScopedLock lock(streamSetLock); - streamSet.insert(this); - } + amd::ScopedLock lock(streamSetLock); + streamSet.insert(this); } else { Destroy(); } @@ -104,6 +102,15 @@ int Stream::DeviceId() const { return device_->deviceId(); } +void Stream::syncNonBlockingStreams() { + amd::ScopedLock lock(streamSetLock); + for (auto& it : streamSet) { + if (it->Flags() & hipStreamNonBlocking) { + it->asHostQueue()->finish(); + } + } +} + }; // ================================================================================================ @@ -116,6 +123,8 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream amd::HostQueue* active_queue = stream->asHostQueue(); // If it's the current device if ((&active_queue->device() == &blocking_queue->device()) && + // Make sure it's a default stream + ((stream->Flags() & hipStreamNonBlocking) == 0) && // and it's not the current stream (active_queue != blocking_queue) && // check for a wait on the null stream @@ -199,7 +208,7 @@ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, priority = static_cast(amd::CommandQueue::Priority::Normal); } - return HIP_RETURN(ihipStreamCreate(stream, flags, static_cast(priority))); + HIP_RETURN(ihipStreamCreate(stream, flags, static_cast(priority))); } // ================================================================================================ @@ -213,7 +222,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio // Only report one kind of priority for now. *greatestPriority = static_cast(amd::CommandQueue::Priority::Normal); } - return HIP_RETURN(hipSuccess); + HIP_RETURN(hipSuccess); } // ================================================================================================ @@ -330,3 +339,15 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv)); } +// ================================================================================================ +hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { + HIP_INIT_API(hipStreamGetPriority, stream, priority); + if ((priority != nullptr) && (stream != nullptr)) { + *priority = static_cast(reinterpret_cast(stream)->Priority()); + } else { + HIP_RETURN(hipErrorInvalidValue); + } + + HIP_RETURN(hipSuccess); + +} diff --git a/rocclr/hip_texture.cpp b/rocclr/hip_texture.cpp index a9121f183b..9d16e3da01 100755 --- a/rocclr/hip_texture.cpp +++ b/rocclr/hip_texture.cpp @@ -772,7 +772,7 @@ hipError_t hipTexRefGetArray(hipArray_t* pArray, // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace. hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef->textureObject); if (error != hipSuccess) { - return HIP_RETURN(error); + HIP_RETURN(error); } switch (resDesc.resType) { @@ -848,7 +848,7 @@ hipError_t hipTexRefGetAddress(hipDeviceptr_t* dptr, if (error != hipSuccess) { DevLogPrintfError("hipGetTextureObjectResourceDesc failed with error code: %s \n", hipGetErrorName(error)); - return HIP_RETURN(error); + HIP_RETURN(error); } switch (resDesc.resType) { @@ -903,7 +903,7 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset, // Align the user ptr to HW requirments. resDesc.res.linear.devPtr = static_cast(dptr) - *ByteOffset; } else { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidValue); } hipTextureDesc texDesc = hip::getTextureDesc(texRef); @@ -1108,7 +1108,7 @@ hipError_t hipTexRefGetMipmappedArray(hipMipmappedArray_t* pArray, // TODO use ihipGetTextureObjectResourceDesc() to not pollute the API trace. hipError_t error = hipGetTextureObjectResourceDesc(&resDesc, texRef->textureObject); if (error != hipSuccess) { - return HIP_RETURN(error); + HIP_RETURN(error); } switch (resDesc.resType) { diff --git a/tests/src/runtimeApi/stream/hipMultiStreams.cpp b/tests/src/runtimeApi/stream/hipMultiStreams.cpp new file mode 100644 index 0000000000..a6cc18f6e9 --- /dev/null +++ b/tests/src/runtimeApi/stream/hipMultiStreams.cpp @@ -0,0 +1,89 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +using namespace std; + +__global__ void kernel_do_nothing() { + // empty kernel +} + +int main(int argc, char* argv[]) { + + constexpr int nLoops = 100000; + constexpr int nStreams = 2; + vector streams(nStreams); + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return 0; + } + + static int device = 0; + HIPCHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << endl; + + for (int i = 0; i < nStreams; i++) { + HIPCHECK(hipStreamCreate(&streams[i])); + } + + for (int k = 0; k <= nLoops; ++k) { + HIPCHECK(hipDeviceSynchronize()); + + // Launch kernel with default stream + hipLaunchKernelGGL((kernel_do_nothing), dim3(1), dim3(1), 0, 0); + + // Launch kernel on all streams + for (int i = 0; i < nStreams; i++) { + hipLaunchKernelGGL((kernel_do_nothing), dim3(1), dim3(1), 0, streams[i]); + } + + // Sync stream 1 + HIPCHECK(hipStreamSynchronize(streams[0])); + + if (k % 10000 == 0 || k == nLoops) { + cout << "Info: Iteration = " << k << endl; + } + } + + HIPCHECK(hipDeviceSynchronize()); + + // Clean up + for (int i = 0; i < nStreams; i++) { + HIPCHECK(hipStreamDestroy(streams[i])); + } + + HIPCHECK(hipDeviceReset()); + + passed(); +} diff --git a/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp b/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp new file mode 100644 index 0000000000..94ac8bb34c --- /dev/null +++ b/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp @@ -0,0 +1,164 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include +#include +#include +#include "test_common.h" + +using namespace std; + +__global__ void vector_square(float *C_d, float *A_d, size_t N) { + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + for (size_t i = idx; i < N; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + +int main(int argc, char* argv[]) { + constexpr uint32_t numPartition = 4; + float *dA[numPartition], *dC[numPartition]; + float *hA, *hC; + size_t N = 25 * 1024 * 1024; + size_t Nbytes = N * sizeof(float); + vector streams(numPartition); + vector> cuMasks(numPartition); + stringstream ss[numPartition]; + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + if (nGpu < 1) { + cout << "info: didn't find any GPU! skipping the test!\n"; + passed(); + return 0; + } + + static int device = 0; + HIPCHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << + " with " << props.multiProcessorCount << " CUs" << endl; + + hA = new float[Nbytes]; + HIPCHECK(hA == 0 ? hipErrorOutOfMemory : hipSuccess); + hC = new float[Nbytes]; + HIPCHECK(hC == 0 ? hipErrorOutOfMemory : hipSuccess); + for (size_t i = 0; i < N; i++) { + hA[i] = 1.618f + i; + } + + for (int np = 0; np < numPartition; np++) { + + HIPCHECK(hipMalloc(&dA[np], Nbytes)); + HIPCHECK(hipMalloc(&dC[np], Nbytes)); + + // make unique CU masks in the multiple of dwords for each stream + uint32_t temp = 0; + uint32_t bit_index = np; + for (int i = np; i < props.multiProcessorCount; i = i + 4) { + temp |= 1UL << bit_index; + if (bit_index >= 32) { + cuMasks[np].push_back(temp); + temp = 0; + bit_index = np; + temp |= 1UL << bit_index; + } + bit_index += 4; + } + if (bit_index != 0) { + cuMasks[np].push_back(temp); + } + + HIPCHECK(hipExtStreamCreateWithCUMask(&streams[np], cuMasks[np].size(), cuMasks[np].data())); + + HIPCHECK(hipMemcpy(dA[np], hA, Nbytes, hipMemcpyHostToDevice)); + + ss[np] << std::hex; + for (int i = cuMasks[np].size() - 1; i >= 0; i--) { + ss[np] << cuMasks[np][i]; + } + } + + const unsigned blocks = 512; + const unsigned threadsPerBlock = 256; + + auto single_start = chrono::steady_clock::now(); + cout << "info: launch 'vector_square' kernel on one stream " << streams[0] << " with CU mask: 0x" << ss[0].str().c_str() << endl; + + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, streams[0], dC[0], dA[0], N); + hipDeviceSynchronize(); + + auto single_end = chrono::steady_clock::now(); + chrono::duration single_kernel_time = single_end - single_start; + + HIPCHECK(hipMemcpy(hC, dC[0], Nbytes, hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < N; i++) { + if (hC[i] != hA[i] * hA[i]) { + cout << "info: validation failed for kernel launched at stream" << streams[0] << endl; + HIPCHECK(hipErrorUnknown); + } + } + + cout << "info: launch 'vector_square' kernel on " << numPartition << " streams:" << endl; + auto all_start = chrono::steady_clock::now(); + for (int np = 0; np < numPartition; np++) { + cout << "info: launch 'vector_square' kernel on the stream " << streams[np] << " with CU mask: 0x" << ss[np].str().c_str() << endl; + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + streams[np], dC[np], dA[np], N); + } + hipDeviceSynchronize(); + + auto all_end = chrono::steady_clock::now(); + chrono::duration all_kernel_time = all_end - all_start; + + for (int np = 0; np < numPartition; np++) { + HIPCHECK(hipMemcpy(hC, dC[np], Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < N; i++) { + if (hC[i] != hA[i] * hA[i]) { + cout << "info: validation failed for kernel launched at stream" << streams[np] << endl; + HIPCHECK(hipErrorUnknown); + } + } + } + + cout << "info: kernel launched on one stream took: " << single_kernel_time.count() << " seconds" << endl; + cout << "info: kernels launched on " << numPartition << " streams took: " << all_kernel_time.count() << " seconds" << endl; + cout << "info: launching kernels on " << numPartition << " streams asynchronously is " << single_kernel_time.count() / (all_kernel_time.count() / numPartition) + << " times faster per stream than launching on one stream alone" << endl; + + delete [] hA; + delete [] hC; + for (int np = 0; np < numPartition; np++) { + hipFree(dC[np]); + hipFree(dA[np]); + HIPCHECK(hipStreamDestroy(streams[np])); + } + + passed(); +}