SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: Iba0157eecbc557030c3b77e181cd0fab171cdf4c
This commit is contained in:
Jenkins
Vendored
+50
-5
@@ -28,7 +28,43 @@ def hipBuildTest(String backendLabel) {
|
||||
}
|
||||
}
|
||||
}
|
||||
stage("build") {
|
||||
stage("Build - HIT framework") {
|
||||
// Running the build on hipamd workspace
|
||||
dir("${WORKSPACE}/hipamd") {
|
||||
sh """#!/usr/bin/env bash
|
||||
set -x
|
||||
rm -rf build
|
||||
mkdir build
|
||||
cd build
|
||||
# Check if backend label contains string "amd" or backend host is a server with amd gpu
|
||||
if [[ $backendLabel =~ amd ]]; then
|
||||
cmake -DHIP_CATCH_TEST=0 -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=\$OPENCL_DIR -DROCCLR_PATH=\$ROCclr_DIR -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
else
|
||||
cmake -DHIP_CATCH_TEST=0 -DHIP_PLATFORM=nvidia -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
fi
|
||||
make install -j\$(nproc)
|
||||
if [[ $backendLabel =~ amd ]]; then
|
||||
make build_tests -j\$(nproc)
|
||||
else
|
||||
HIP_COMPILER=nvcc HIP_PLATFORM=nvidia make build_tests -j\$(nproc)
|
||||
fi
|
||||
"""
|
||||
}
|
||||
}
|
||||
stage('HIP Unit Tests - HIT framework') {
|
||||
dir("${WORKSPACE}/hipamd/build") {
|
||||
sh """#!/usr/bin/env bash
|
||||
set -x
|
||||
# Check if backend label contains string "amd" or backend host is a server with amd gpu
|
||||
if [[ $backendLabel =~ amd ]]; then
|
||||
LLVM_PATH=/opt/rocm/llvm make test
|
||||
else
|
||||
make test
|
||||
fi
|
||||
"""
|
||||
}
|
||||
}
|
||||
stage("Build - Catch2 framework") {
|
||||
// Running the build on hipamd workspace
|
||||
dir("${WORKSPACE}/hipamd") {
|
||||
sh """#!/usr/bin/env bash
|
||||
@@ -37,16 +73,25 @@ def hipBuildTest(String backendLabel) {
|
||||
cd build
|
||||
# Check if backend label contains string "amd" or backend host is a server with amd gpu
|
||||
if [[ $backendLabel =~ amd ]]; then
|
||||
cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=\$OPENCL_DIR -DROCCLR_PATH=\$ROCclr_DIR -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
cmake -DHIP_CATCH_TEST=1 -DHIP_PATH=\$PWD/install -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=\$OPENCL_DIR -DROCCLR_PATH=\$ROCclr_DIR -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
else
|
||||
cmake -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
export HIP_PLATFORM=nvidia
|
||||
export HIP_COMPILER=nvcc
|
||||
export HIP_RUNTIME=cuda
|
||||
cmake -DHIP_CATCH_TEST=1 -DHIP_PATH=\$PWD/install -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_INSTALL_PREFIX=\$PWD/install ..
|
||||
fi
|
||||
|
||||
make install -j\$(nproc)
|
||||
make build_tests -j\$(nproc)
|
||||
|
||||
if [[ $backendLabel =~ amd ]]; then
|
||||
make build_tests -j\$(nproc)
|
||||
else
|
||||
HIP_COMPILER=nvcc HIP_PLATFORM=nvidia make build_tests -j\$(nproc)
|
||||
fi
|
||||
"""
|
||||
}
|
||||
}
|
||||
stage('HIP Unit Tests') {
|
||||
stage('HIP Unit Tests - Catch2 framework') {
|
||||
dir("${WORKSPACE}/hipamd/build") {
|
||||
sh """#!/usr/bin/env bash
|
||||
set -x
|
||||
|
||||
+7
-6
@@ -1,7 +1,8 @@
|
||||
# Contributor Guidelines
|
||||
|
||||
## Make Tips
|
||||
When building HIP, you will likely want to build and install to a local user-accessible directory (rather than /opt/rocm).
|
||||
ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is /opt/rocm.
|
||||
When building HIP, you will likely want to build and install to a local user-accessible directory (rather than <ROCM_PATH>).
|
||||
This can be easily be done by setting the -DCMAKE_INSTALL_PREFIX variable when running cmake. Typical use case is to
|
||||
set CMAKE_INSTALL_PREFIX to your HIP git root, and then ensure HIP_PATH points to this directory. For example
|
||||
|
||||
@@ -31,7 +32,7 @@ In some cases new HIP-Clang features are tied to specified releases, and it can
|
||||
HIP runtime version
|
||||
|
||||
```
|
||||
> cat /opt/rocm/hip/bin/.hipVersion
|
||||
> cat <ROCM_PATH>/hip/bin/.hipVersion
|
||||
# Auto-generated by cmake
|
||||
HIP_VERSION_MAJOR=3
|
||||
HIP_VERSION_MINOR=9
|
||||
@@ -41,11 +42,11 @@ HIP_VERSION_PATCH=20345-519ef3f2
|
||||
HIP-Clang compiler version
|
||||
|
||||
```
|
||||
$ /opt/rocm/llvm/bin/clang -v
|
||||
$ <ROCM_PATH>/llvm/bin/clang -v
|
||||
clang version 11.0.0 (/src/external/llvm-project/clang 075fedd3fd2f4d9d8cca79d0cd51f64c5ef21432)
|
||||
Target: x86_64-unknown-linux-gnu
|
||||
Thread model: posix
|
||||
InstalledDir: /opt/rocm/llvm/bin
|
||||
InstalledDir: <ROCM_PATH>/llvm/bin
|
||||
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/7
|
||||
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/7.5.0
|
||||
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/8
|
||||
@@ -73,8 +74,8 @@ For applications and benchmarks outside the directed test environment, developme
|
||||
|
||||
## Environment Variables
|
||||
- **HIP_PATH** : Location of HIP include, src, bin, lib directories.
|
||||
- **HCC_ROCCLR_HOME** : Path to HIP/ROCclr directory, used on AMD platforms. Default /opt/rocm/rocclr.
|
||||
- **HSA_PATH** : Path to HSA include, lib. Default /opt/rocm/hsa.
|
||||
- **HCC_ROCCLR_HOME** : Path to HIP/ROCclr directory, used on AMD platforms. Default <ROCM_PATH>/rocclr.
|
||||
- **HSA_PATH** : Path to HSA include, lib. Default <ROCM_PATH>/hsa.
|
||||
- **CUDA_PATH* : On nvcc system, this points to root of CUDA installation.
|
||||
|
||||
## Contribution guidelines ##
|
||||
|
||||
+14
-14
@@ -21,7 +21,7 @@ HIP can be easily installed using pre-built binary packages using the package ma
|
||||
HIP code can be developed either on AMD ROCm platform using HIP-Clang compiler, or a CUDA platform with nvcc installed.
|
||||
|
||||
## AMD Platform
|
||||
|
||||
ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is /opt/rocm.
|
||||
```
|
||||
sudo apt install mesa-common-dev
|
||||
sudo apt install clang
|
||||
@@ -38,17 +38,17 @@ HIP-Clang can be built manually:
|
||||
git clone -b amd-stg-open https://github.com/RadeonOpenCompute/llvm-project.git
|
||||
cd llvm-project
|
||||
mkdir -p build && cd build
|
||||
cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=1 -DLLVM_TARGETS_TO_BUILD="AMDGPU;X86" -DLLVM_ENABLE_PROJECTS="clang;lld;compiler-rt" ../llvm
|
||||
cmake -DCMAKE_INSTALL_PREFIX=<ROCM_PATH>/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=1 -DLLVM_TARGETS_TO_BUILD="AMDGPU;X86" -DLLVM_ENABLE_PROJECTS="clang;lld;compiler-rt" ../llvm
|
||||
make -j
|
||||
sudo make install
|
||||
```
|
||||
Rocm device library can be manually built as following,
|
||||
```
|
||||
export PATH=/opt/rocm/llvm/bin:$PATH
|
||||
export PATH=<ROCM_PATH>/llvm/bin:$PATH
|
||||
git clone -b amd-stg-open https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git
|
||||
cd ROCm-Device-Libs
|
||||
mkdir -p build && cd build
|
||||
CC=clang CXX=clang++ cmake -DLLVM_DIR=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_WERROR=1 -DLLVM_ENABLE_ASSERTIONS=1 -DCMAKE_INSTALL_PREFIX=/opt/rocm ..
|
||||
CC=clang CXX=clang++ cmake -DLLVM_DIR=<ROCM_PATH>/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_WERROR=1 -DLLVM_ENABLE_ASSERTIONS=1 -DCMAKE_INSTALL_PREFIX=<ROCM_PATH> ..
|
||||
make -j
|
||||
sudo make install
|
||||
```
|
||||
@@ -64,8 +64,8 @@ apt-get install hip-runtime-nvidia hip-devel
|
||||
```
|
||||
* Default paths and environment variables:
|
||||
* By default HIP looks for CUDA SDK in /usr/local/cuda (can be overriden by setting CUDA_PATH env variable).
|
||||
* By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
|
||||
* By default HIP is installed into <ROCM_PATH>/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* Optionally, consider adding <ROCM_PATH>/bin to your path to make it easier to use the tools.
|
||||
|
||||
# Building HIP from source
|
||||
|
||||
@@ -98,21 +98,21 @@ See https://github.com/ROCm-Developer-Tools/hipamd
|
||||
```
|
||||
cd "$HIPAMD_DIR"
|
||||
mkdir -p build; cd build
|
||||
cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=$OPENCL_DIR -DROCCLR_PATH=$ROCCLR_DIR -DCMAKE_PREFIX_PATH="/opt/rocm/" -DCMAKE_INSTALL_PREFIX=$PWD/install ..
|
||||
cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=$OPENCL_DIR -DROCCLR_PATH=$ROCCLR_DIR -DCMAKE_PREFIX_PATH="<ROCM_PATH>/" -DCMAKE_INSTALL_PREFIX=$PWD/install ..
|
||||
make -j$(nproc)
|
||||
sudo make install
|
||||
```
|
||||
|
||||
Note: If you don't specify CMAKE_INSTALL_PREFIX, hip runtime will be installed to "/opt/rocm/hip".
|
||||
Note: If you don't specify CMAKE_INSTALL_PREFIX, hip runtime will be installed to "<ROCM_PATH>/hip".
|
||||
By default, release version of AMDHIP is built.
|
||||
|
||||
## Default paths and environment variables
|
||||
|
||||
* By default HIP looks for HSA in /opt/rocm/hsa (can be overridden by setting HSA_PATH environment variable).
|
||||
* By default HIP is installed into /opt/rocm/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* By default HIP looks for clang in /opt/rocm/llvm/bin (can be overridden by setting HIP_CLANG_PATH environment variable)
|
||||
* By default HIP looks for device library in /opt/rocm/lib (can be overridden by setting DEVICE_LIB_PATH environment variable).
|
||||
* Optionally, consider adding /opt/rocm/bin to your PATH to make it easier to use the tools.
|
||||
* By default HIP looks for HSA in <ROCM_PATH>/hsa (can be overridden by setting HSA_PATH environment variable).
|
||||
* By default HIP is installed into <ROCM_PATH>/hip (can be overridden by setting HIP_PATH environment variable).
|
||||
* By default HIP looks for clang in <ROCM_PATH>/llvm/bin (can be overridden by setting HIP_CLANG_PATH environment variable)
|
||||
* By default HIP looks for device library in <ROCM_PATH>/lib (can be overridden by setting DEVICE_LIB_PATH environment variable).
|
||||
* Optionally, consider adding <ROCM_PATH>/bin to your PATH to make it easier to use the tools.
|
||||
* Optionally, set HIPCC_VERBOSE=7 to output the command line for compilation.
|
||||
|
||||
After installation, make sure HIP_PATH is pointed to /where/to/install/hip
|
||||
@@ -121,7 +121,7 @@ After installation, make sure HIP_PATH is pointed to /where/to/install/hip
|
||||
|
||||
Run hipconfig (instructions below assume default installation path) :
|
||||
```shell
|
||||
/opt/rocm/bin/hipconfig --full
|
||||
<ROCM_PATH>/bin/hipconfig --full
|
||||
```
|
||||
|
||||
Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/main/samples/0_Intro/square).
|
||||
|
||||
@@ -14,9 +14,9 @@ New projects can be developed directly in the portable HIP C++ language and can
|
||||
|
||||
## DISCLAIMER
|
||||
|
||||
The information contained herein is for informational purposes only, and is subject to change without notice. In addition, any stated support is planned and is also subject to change. While every precaution has been taken in the preparation of this document, it may contain technical inaccuracies, omissions and typographical errors, and AMD is under no obligation to update or otherwise correct this information. Advanced Micro Devices, Inc. makes no representations or warranties with respect to the accuracy or completeness of the contents of this document, and assumes no liability of any kind, including the implied warranties of noninfringement, merchantability or fitness for particular purposes, with respect to the operation or use of AMD hardware, software or other products described herein. No license, including implied or arising by estoppel, to any intellectual property rights is granted by this document. Terms and limitations applicable to the purchase or use of AMD’s products are as set forth in a signed agreement between the parties or in AMD's Standard Terms and Conditions of Sale.
|
||||
The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard versionchanges, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated.AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes.THIS INFORMATION IS PROVIDED ‘AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
|
||||
|
||||
© 2020 Advanced Micro Devices, Inc. All Rights Reserved.
|
||||
© 2021 Advanced Micro Devices, Inc. All Rights Reserved.
|
||||
|
||||
## Repository branches:
|
||||
|
||||
@@ -44,6 +44,7 @@ HIP releases are typically naming convention for each ROCM release to help diffe
|
||||
- [HIP Logging ](docs/markdown/hip_logging.md)
|
||||
- [HIP Debugging ](docs/markdown/hip_debugging.md)
|
||||
- [Code Object tooling ](docs/markdown/obj_tooling.md)
|
||||
- [HIP RTC](docs/markdown/hip_rtc.md)
|
||||
- [HIP Terminology](docs/markdown/hip_terms2.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/OpenCL)
|
||||
- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md)
|
||||
- Supported CUDA APIs:
|
||||
@@ -145,11 +146,8 @@ The README with the procedures and tips the team used during this porting effort
|
||||
* **amd_detail/**** , **nvidia_detail/**** : Implementation details for specific platforms. HIP applications should not include these files directly.
|
||||
|
||||
* **bin**: Tools and scripts to help with hip porting
|
||||
* **hipify-perl** : Script based tool to convert CUDA code to portable CPP. Converts CUDA APIs and kernel builtins.
|
||||
* **hipcc** : Compiler driver that can be used to replace nvcc in existing CUDA code. hipcc will call nvcc or HIP-Clang depending on platform and include appropriate platform-specific headers and libraries.
|
||||
* **hipconfig** : Print HIP configuration (HIP_PATH, HIP_PLATFORM, HIP_COMPILER, HIP_RUNTIME, CXX config flags, etc.)
|
||||
* **hipexamine-perl.sh** : Script to scan the directory, find all code, and report statistics on how much can be ported with HIP (and identify likely features not yet supported).
|
||||
* **hipconvertinplace-perl.sh** : Script to scan the directory, find all code, and convert the found CUDA code to HIP reporting all unconverted things.
|
||||
|
||||
* **doc**: Documentation - markdown and doxygen info.
|
||||
|
||||
|
||||
@@ -1,27 +0,0 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
|
||||
SEARCH_DIRS=$@
|
||||
|
||||
find $SEARCH_DIRS -name '*.cu'
|
||||
find $SEARCH_DIRS -name '*.cpp' -o -name '*.cxx' -o -name '*.c' -o -name '*.cc'
|
||||
find $SEARCH_DIRS -name '*.cuh'
|
||||
find $SEARCH_DIRS -name '*.h' -o -name '*.hpp' -o -name '*.inc' -o -name '*.inl' -o -name '*.hxx' -o -name '*.hdl'
|
||||
@@ -1,24 +0,0 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
|
||||
SEARCH_DIR=$1
|
||||
|
||||
find $SEARCH_DIR -not -name '*.cu' -and -not -name '*.cpp' -and -not -name '*.cxx' -and -not -name '*.c' -and -not -name '*.cc' -and -not -name '*.cuh' -and -not -name '*.h' -and -not -name '*.hpp' -and -not -name '*.inc' -and -not -name '*.inl' -and -not -name '*.hxx' -and -not -name '*.hdl'
|
||||
@@ -58,11 +58,6 @@ $verbose = $ENV{'HIPCC_VERBOSE'} // 0;
|
||||
$HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'};
|
||||
$HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'};
|
||||
|
||||
# Known HIP target names.
|
||||
@knownTargets = ('gfx700', 'gfx701', 'gfx702', 'gfx703', 'gfx704', 'gfx705',
|
||||
'gfx801', 'gfx802', 'gfx803', 'gfx805', 'gfx810',
|
||||
'gfx900', 'gfx902', 'gfx904', 'gfx906', 'gfx908', 'gfx909', 'gfx90a',
|
||||
'gfx1010', 'gfx1011', 'gfx1012', 'gfx1030', 'gfx1031', 'gfx1032');
|
||||
# Known Features
|
||||
@knownFeatures = ('sramecc-', 'sramecc+', 'xnack-', 'xnack+');
|
||||
|
||||
@@ -146,11 +141,15 @@ $setStdLib = 0; # TODO - set to 0
|
||||
$default_amdgpu_target = 1;
|
||||
|
||||
if ($HIP_PLATFORM eq "amd") {
|
||||
$HIPCC="\"$HIP_CLANG_PATH/clang++\"";
|
||||
$execExtension = "";
|
||||
if($isWindows) {
|
||||
$execExtension = ".exe";
|
||||
}
|
||||
$HIPCC="$HIP_CLANG_PATH/clang++" . $execExtension;
|
||||
|
||||
# If $HIPCC clang++ is not compiled, use clang instead
|
||||
if ( ! -e $HIPCC ) {
|
||||
$HIPCC="\"$HIP_CLANG_PATH/clang\"";
|
||||
$HIPCC="$HIP_CLANG_PATH/clang" . $execExtension;
|
||||
$HIPLDFLAGS = "--driver-mode=g++";
|
||||
}
|
||||
|
||||
@@ -158,6 +157,9 @@ if ($HIP_PLATFORM eq "amd") {
|
||||
$HIP_CLANG_VERSION=~/.*clang version (\S+).*/;
|
||||
$HIP_CLANG_VERSION=$1;
|
||||
|
||||
# Figure out the target with which llvm is configured
|
||||
$HIP_CLANG_TARGET = `$HIPCC -print-target-triple`;
|
||||
|
||||
if (! defined $HIP_CLANG_INCLUDE_PATH) {
|
||||
$HIP_CLANG_INCLUDE_PATH = abs_path("$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include");
|
||||
}
|
||||
@@ -177,6 +179,7 @@ if ($HIP_PLATFORM eq "amd") {
|
||||
print ("HIP_INCLUDE_PATH=$HIP_INCLUDE_PATH\n");
|
||||
print ("HIP_LIB_PATH=$HIP_LIB_PATH\n");
|
||||
print ("DEVICE_LIB_PATH=$DEVICE_LIB_PATH\n");
|
||||
print ("HIP_CLANG_TARGET=$HIP_CLANG_TARGET\n");
|
||||
}
|
||||
|
||||
if ($isWindows) {
|
||||
@@ -633,11 +636,6 @@ if($HIP_PLATFORM eq "amd"){
|
||||
if ($HIP_PLATFORM eq 'amd' and $hasHIP) {
|
||||
$HIPCXXFLAGS .= $GPU_ARCH_ARG;
|
||||
}
|
||||
|
||||
# If the specified target is not in the list of known target names, emit a warning.
|
||||
if (grep($procName, @knownTargets) eq 0) {
|
||||
print "Warning: The specified HIP target: $val is unknown. Correct compilation is not guaranteed.\n";
|
||||
}
|
||||
}
|
||||
}
|
||||
if ($hsacoVersion > 0) {
|
||||
@@ -711,7 +709,12 @@ if ($HIP_PLATFORM eq "amd") {
|
||||
$toolArgs = " -Wl,--enable-new-dtags -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 " . ${toolArgs};
|
||||
}
|
||||
# To support __fp16 and _Float16, explicitly link with compiler-rt
|
||||
$toolArgs .= " -L$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/linux -lclang_rt.builtins-x86_64 "
|
||||
$HIP_CLANG_BUILTIN_LIB="$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/$HIP_CLANG_TARGET/libclang_rt.builtins.a";
|
||||
if (-e $HIP_CLANG_BUILTIN_LIB) {
|
||||
$toolArgs .= " -L$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/$HIP_CLANG_TARGET -lclang_rt.builtins "
|
||||
} else {
|
||||
$toolArgs .= " -L$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/linux -lclang_rt.builtins-x86_64 "
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
+2
-2
@@ -201,7 +201,7 @@ if (!$printed or $p_full) {
|
||||
print "=== Environment Variables\n";
|
||||
if ($isWindows) {
|
||||
print ("PATH=$ENV{PATH}\n");
|
||||
system("set | findstr /B /C:\"HIP\" /C:\"HSA\" /C:\"CUDA\" /C:\"LD_LIBRARY_PATH\"");
|
||||
system("set | findstr //B //C:\"HIP\" //C:\"HSA\" //C:\"CUDA\" //C:\"LD_LIBRARY_PATH\"");
|
||||
} else {
|
||||
system("echo PATH=\$PATH");
|
||||
system("env | egrep '^HIP|^HSA|^CUDA|^LD_LIBRARY_PATH'");
|
||||
@@ -212,7 +212,7 @@ if (!$printed or $p_full) {
|
||||
if ($isWindows) {
|
||||
print "== Windows Display Drivers\n";
|
||||
print "Hostname : "; system ("hostname");
|
||||
system ("wmic path win32_VideoController get AdapterCompatibility,InstalledDisplayDrivers,Name | findstr /B /C:\"Advanced Micro Devices\"");
|
||||
system ("wmic path win32_VideoController get AdapterCompatibility,InstalledDisplayDrivers,Name | findstr //B //C:\"Advanced Micro Devices\"");
|
||||
} else {
|
||||
print "== Linux Kernel\n";
|
||||
print "Hostname : "; system ("hostname");
|
||||
|
||||
@@ -1,37 +0,0 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2017 - 2021 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.
|
||||
|
||||
#usage : hipconvertinplace-perl.sh DIRNAME [hipify-perl options]
|
||||
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
# This can be quite handy when dealing with an existing CUDA code base since the script
|
||||
# preserves the existing directory structure.
|
||||
|
||||
# For each code file, this script will:
|
||||
# - If ".prehip file does not exist, copy the original code to a new file with extension ".prehip". Then hipify the code file.
|
||||
# - If ".prehip" file exists, this is used as input to hipify.
|
||||
# (this is useful for testing improvements to the hipify-perl toolset).
|
||||
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify-perl -inplace -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
@@ -1,43 +0,0 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
|
||||
#usage : hipconvertinplace.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
#hipify "inplace" all code files in specified directory.
|
||||
# This can be quite handy when dealing with an existing CUDA code base since the script
|
||||
# preserves the existing directory structure.
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -inplace -print-stats $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
@@ -1,41 +0,0 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
|
||||
#usage : hipexamine.sh DIRNAME [hipify options] [--] [clang options]
|
||||
|
||||
# Generate CUDA->HIP conversion statistics for all the code files in the specified directory.
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
|
||||
hipify_args=''
|
||||
while (( "$#" )); do
|
||||
shift
|
||||
if [ "$1" != "--" ]; then
|
||||
hipify_args="$hipify_args $1"
|
||||
else
|
||||
shift
|
||||
break
|
||||
fi
|
||||
done
|
||||
clang_args="$@"
|
||||
|
||||
$SCRIPT_DIR/hipify-clang -examine $hipify_args `$SCRIPT_DIR/findcode.sh $SEARCH_DIR` -- -x cuda $clang_args
|
||||
@@ -1,280 +0,0 @@
|
||||
#!/usr/bin/env perl
|
||||
##
|
||||
# Copyright (c) 2015 - 2021 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.
|
||||
##
|
||||
#usage hipify-cmakefile [OPTIONS] INPUT_FILE
|
||||
use Getopt::Long;
|
||||
use warnings;
|
||||
|
||||
GetOptions(
|
||||
"print-stats" => \$print_stats # print the command-line, like a header.
|
||||
, "quiet-warnings" => \$quiet_warnings # don't print warnings on unknown CUDA functions.
|
||||
, "no-output" => \$no_output # don't write any translated output to stdout.
|
||||
, "inplace" => \$inplace # modify input file inplace, save backup in ".prehip" file.
|
||||
, "n" => \$n # combination of print_stats + no-output.
|
||||
);
|
||||
|
||||
$print_stats = 1 if $n;
|
||||
$no_output = 1 if $n;
|
||||
|
||||
@warn_whitelist = ();
|
||||
|
||||
#---
|
||||
#Stats tracking code:
|
||||
@statNames = ( "macro", "include", "option", "other" );
|
||||
|
||||
#---
|
||||
#Compute total of all individual counts:
|
||||
sub totalStats {
|
||||
my %count = %{ shift() };
|
||||
|
||||
my $total = 0;
|
||||
foreach $key ( keys %count ) {
|
||||
$total += $count{$key};
|
||||
}
|
||||
|
||||
return $total;
|
||||
}
|
||||
|
||||
#---
|
||||
sub printStats {
|
||||
my $label = shift();
|
||||
my @statNames = @{ shift() };
|
||||
my %counts = %{ shift() };
|
||||
my $warnings = shift();
|
||||
my $loc = shift();
|
||||
|
||||
my $total = totalStats( \%counts );
|
||||
|
||||
printf STDERR "%s %d CUDA->HIP refs( ", $label, $total;
|
||||
|
||||
foreach $stat (@statNames) {
|
||||
printf STDERR "%s:%d ", $stat, $counts{$stat};
|
||||
}
|
||||
|
||||
printf STDERR ") warn:%d LOC:%d", $warnings, $loc;
|
||||
}
|
||||
|
||||
#---
|
||||
# Add adder stats to dest. Used to add stats for current file to a running total for all files:
|
||||
sub addStats {
|
||||
my $dest_ref = shift();
|
||||
my %adder = %{ shift() };
|
||||
|
||||
foreach $key ( keys %adder ) {
|
||||
$dest_ref->{$key} += $adder{$key};
|
||||
}
|
||||
}
|
||||
|
||||
#---
|
||||
sub clearStats {
|
||||
my $dest_ref = shift();
|
||||
my @statNames = @{ shift() };
|
||||
|
||||
foreach $stat (@statNames) {
|
||||
$dest_ref->{$stat} = 0;
|
||||
}
|
||||
}
|
||||
|
||||
# count of transforms in all files:
|
||||
my %tt;
|
||||
clearStats( \%tt, \@statNames );
|
||||
|
||||
my $fileCount = @ARGV;
|
||||
my $fileName = "";
|
||||
|
||||
while (@ARGV) {
|
||||
$fileName = shift(@ARGV);
|
||||
if ($inplace) {
|
||||
my $file_prehip = "$fileName" . ".prehip";
|
||||
my $infile;
|
||||
my $outfile;
|
||||
if ( -e $file_prehip ) {
|
||||
$infile = $file_prehip;
|
||||
$outfile = $fileName;
|
||||
}
|
||||
else {
|
||||
system("cp $fileName $file_prehip");
|
||||
$infile = $file_prehip;
|
||||
$outfile = $fileName;
|
||||
}
|
||||
open( INFILE, "<", $infile ) or die "error: could not open $infile";
|
||||
open( OUTFILE, ">", $outfile ) or die "error: could not open $outfile";
|
||||
$OUTFILE = OUTFILE;
|
||||
}
|
||||
else {
|
||||
open( INFILE, "<", $fileName ) or die "error: could not open $fileName";
|
||||
$OUTFILE = STDOUT;
|
||||
}
|
||||
|
||||
# count of transforms in this file, init to 0 here:
|
||||
my %ft;
|
||||
clearStats( \%ft, \@statNames );
|
||||
|
||||
my $lineCount = 0;
|
||||
|
||||
undef $/; # Read whole file at once, so we can match newlines.
|
||||
while (<INFILE>) {
|
||||
|
||||
# Replace find_package(CUDA) with find_package(HIP)
|
||||
$ft{'include'} += s/\bfind_package[ ]*\([ ]*CUDA[ ]*[0-9.]*/find_package(HIP/ig;
|
||||
|
||||
# Replace macros
|
||||
$ft{'macro'} += s/\bCUDA_ADD_EXECUTABLE/HIP_ADD_EXECUTABLE/ig;
|
||||
$ft{'macro'} += s/\bCUDA_ADD_LIBRARY/HIP_ADD_LIBRARY/ig;
|
||||
$ft{'macro'} += s/\bCUDA_INCLUDE_DIRECTORIES/HIP_INCLUDE_DIRECTORIES/ig;
|
||||
|
||||
# Replace options
|
||||
$ft{'option'} += s/\bCUDA_NVCC_FLAGS/HIP_NVCC_FLAGS/ig;
|
||||
$ft{'option'} += s/\bCUDA_HOST_COMPILATION_CPP/HIP_HOST_COMPILATION_CPP/ig;
|
||||
$ft{'option'} += s/\bCUDA_SOURCE_PROPERTY_FORMAT/HIP_SOURCE_PROPERTY_FORMAT/ig;
|
||||
|
||||
# Replace variables
|
||||
$ft{'other'} += s/\bCUDA_FOUND/HIP_FOUND/ig;
|
||||
$ft{'other'} += s/\bCUDA_VERSION/HIP_VERSION/ig;
|
||||
$ft{'other'} += s/\bCUDA_TOOLKIT_ROOT_DIR/HIP_ROOT_DIR/ig;
|
||||
|
||||
unless ($quiet_warnings) {
|
||||
|
||||
#print STDERR "Check WARNINGs\n";
|
||||
# copy into array of lines, process line-by-line to show warnings:
|
||||
my @lines = split /\n/, $_;
|
||||
my $tmp = $_; # copies the whole file, could be a little smarter here...
|
||||
my $line_num = 0;
|
||||
|
||||
foreach (@lines) {
|
||||
$line_num++;
|
||||
|
||||
# remove any whitelisted words:
|
||||
foreach $w (@warn_whitelist) {
|
||||
s/\b$w\b/ZAP/;
|
||||
}
|
||||
|
||||
$s = warnUnsupportedSpecialFunctions($line_num);
|
||||
$warnings += $s;
|
||||
}
|
||||
|
||||
$_ = $tmp;
|
||||
}
|
||||
|
||||
#--------
|
||||
# Print it!
|
||||
unless ($no_output) {
|
||||
print $OUTFILE "$_";
|
||||
}
|
||||
$lineCount = $_ =~ tr/\n//;
|
||||
}
|
||||
|
||||
my $totalConverted = totalStats( \%ft );
|
||||
|
||||
if ( ( $totalConverted + $warnings ) and $print_stats ) {
|
||||
printStats( "info: converted", \@statNames, \%ft, $warnings, $lineCount );
|
||||
print STDERR " in '$fileName'\n";
|
||||
print STDERR "You may need to hand-edit '$fileName' to add steps to build correctly on HCC path\n";
|
||||
}
|
||||
|
||||
# Update totals for all files:
|
||||
addStats( \%tt, \%ft );
|
||||
$Twarnings += $warnings;
|
||||
$TlineCount += $lineCount;
|
||||
}
|
||||
|
||||
#-- Print total stats for all files processed:
|
||||
if ( $print_stats and ( $fileCount > 1 ) ) {
|
||||
print STDERR "\n";
|
||||
printStats( "info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount );
|
||||
print STDERR "\n";
|
||||
}
|
||||
|
||||
#---
|
||||
sub warnUnsupportedSpecialFunctions {
|
||||
my $line_num = shift;
|
||||
my $m = 0;
|
||||
|
||||
foreach $func (
|
||||
# macros:
|
||||
"CUDA_ADD_CUFFT_TO_TARGET",
|
||||
"CUDA_ADD_CUBLAS_TO_TARGET",
|
||||
#"CUDA_ADD_EXECUTABLE",
|
||||
#"CUDA_ADD_LIBRARY",
|
||||
"CUDA_BUILD_CLEAN_TARGET",
|
||||
"CUDA_COMPILE",
|
||||
"CUDA_COMPILE_PTX",
|
||||
"CUDA_COMPILE_FATBIN",
|
||||
"CUDA_COMPILE_CUBIN",
|
||||
"CUDA_COMPUTE_SEPARABLE_COMPILATION_OBJECT_FILE_NAME",
|
||||
#"CUDA_INCLUDE_DIRECTORIES",
|
||||
"CUDA_LINK_SEPARABLE_COMPILATION_OBJECTS",
|
||||
"CUDA_SELECT_NVCC_ARCH_FLAGS",
|
||||
"CUDA_WRAP_SRCS",
|
||||
|
||||
# options:
|
||||
"CUDA_64_BIT_DEVICE_CODE",
|
||||
"CUDA_ATTACH_VS_BUILD_RULE_TO_CUDA_FILE",
|
||||
"CUDA_BUILD_CUBIN",
|
||||
"CUDA_BUILD_EMULATION",
|
||||
"CUDA_LINK_LIBRARIES_KEYWORD",
|
||||
"CUDA_GENERATED_OUTPUT_DIR",
|
||||
#"CUDA_HOST_COMPILATION_CPP",
|
||||
"CUDA_HOST_COMPILER",
|
||||
#"CUDA_NVCC_FLAGS",
|
||||
#"CUDA_NVCC_FLAGS_<CONFIG>",
|
||||
"CUDA_PROPAGATE_HOST_FLAGS",
|
||||
"CUDA_SEPARABLE_COMPILATION",
|
||||
#"CUDA_SOURCE_PROPERTY_FORMAT",
|
||||
"CUDA_USE_STATIC_CUDA_RUNTIME",
|
||||
"CUDA_VERBOSE_BUILD",
|
||||
|
||||
# others:
|
||||
#"CUDA_VERSION_MAJOR",
|
||||
#"CUDA_VERSION_MINOR",
|
||||
#"CUDA_VERSION",
|
||||
#"CUDA_VERSION_STRING",
|
||||
"CUDA_HAS_FP16",
|
||||
#"CUDA_TOOLKIT_ROOT_DIR",
|
||||
"CUDA_SDK_ROOT_DIR",
|
||||
"CUDA_INCLUDE_DIRS",
|
||||
"CUDA_LIBRARIES",
|
||||
"CUDA_CUFFT_LIBRARIES",
|
||||
"CUDA_CUBLAS_LIBRARIES",
|
||||
"CUDA_cudart_static_LIBRARY",
|
||||
"CUDA_cudadevrt_LIBRARY",
|
||||
"CUDA_cupti_LIBRARY",
|
||||
"CUDA_curand_LIBRARY",
|
||||
"CUDA_cusolver_LIBRARY",
|
||||
"CUDA_cusparse_LIBRARY",
|
||||
"CUDA_npp_LIBRARY",
|
||||
"CUDA_nppc_LIBRARY",
|
||||
"CUDA_nppi_LIBRARY",
|
||||
"CUDA_npps_LIBRARY",
|
||||
"CUDA_nvcuvenc_LIBRARY",
|
||||
"CUDA_nvcuvid_LIBRARY"
|
||||
)
|
||||
{
|
||||
my $mt = m/\b($func)/g;
|
||||
if ($mt) {
|
||||
$m += $mt;
|
||||
print STDERR " warning: $fileName:#$line_num : unsupported macro/option : $_\n";
|
||||
}
|
||||
}
|
||||
|
||||
return $m;
|
||||
}
|
||||
-2694
File diff suppressed because it is too large
Load Diff
+1
-1
@@ -60,7 +60,7 @@ sub can_run {
|
||||
}
|
||||
}
|
||||
|
||||
$isWindows = $^O eq 'MSWin32';
|
||||
$isWindows = ($^O eq 'MSWin32' or $^O eq 'msys');
|
||||
|
||||
#
|
||||
# TODO: Fix rpath LDFLAGS settings
|
||||
|
||||
+149
-64
@@ -13,75 +13,160 @@ The language specification for HIP and CUDA forbid calling a
|
||||
differences in the strictness of this restriction, with HIP exhibiting a tighter
|
||||
adherence to the specification and thus less tolerant of infringing code. The
|
||||
solution is to ensure that all functions which are called in a
|
||||
`__device__` context are correctly annotated to reflect it. An interesting case
|
||||
where these differences emerge is shown below. This relies on a the common
|
||||
[C++ Member Detector idiom][1], as it would be implemented pre C++11):
|
||||
`__device__` context are correctly annotated to reflect it.
|
||||
|
||||
```c++
|
||||
#include <cassert>
|
||||
The following is an example of codes using the specification,
|
||||
```
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <type_traits>
|
||||
#include <random>
|
||||
#include "test_common.h"
|
||||
|
||||
struct aye { bool a[1]; };
|
||||
struct nay { bool a[2]; };
|
||||
static std::random_device dev;
|
||||
static std::mt19937 rng(dev());
|
||||
|
||||
// Dual restriction is necessary in HIP if the detector is to work for
|
||||
// __device__ contexts as well as __host__ ones. NVCC is less strict.
|
||||
template<typename T>
|
||||
__host__ __device__
|
||||
const T& cref_t();
|
||||
|
||||
template<typename T>
|
||||
struct Has_call_operator {
|
||||
// Dual restriction is necessary in HIP if the detector is to work for
|
||||
// __device__ contexts as well as __host__ ones. NVCC is less strict.
|
||||
template<typename C>
|
||||
__host__ __device__
|
||||
static
|
||||
aye test(
|
||||
C const *,
|
||||
typename std::enable_if<
|
||||
(sizeof(cref_t<C>().operator()()) > 0)>::type* = nullptr);
|
||||
static
|
||||
nay test(...);
|
||||
|
||||
enum { value = sizeof(test(static_cast<T*>(0))) == sizeof(aye) };
|
||||
};
|
||||
|
||||
template<typename T, typename U, bool callable = has_call_operator<U>::value>
|
||||
struct Wrapper {
|
||||
template<typename V>
|
||||
V f() const { return T{1}; }
|
||||
};
|
||||
|
||||
|
||||
template<typename T, typename U>
|
||||
struct Wrapper<T, U, true> {
|
||||
template<typename V>
|
||||
V f() const { return T{10}; }
|
||||
};
|
||||
|
||||
// This specialisation will yield a compile-time error, if selected.
|
||||
template<typename T, typename U>
|
||||
struct Wrapper<T, U, false> {};
|
||||
|
||||
template<typename T>
|
||||
struct Functor;
|
||||
|
||||
template<> struct Functor<float> {
|
||||
__device__
|
||||
float operator()() const { return 42.0f; }
|
||||
};
|
||||
|
||||
__device__
|
||||
void this_will_not_compile_if_detector_is_not_marked_device()
|
||||
{
|
||||
float f = Wrapper<float, Functor<float>>().f<float>();
|
||||
template <typename T, typename M>
|
||||
__host__ __device__ inline constexpr int count() {
|
||||
return sizeof(T) / sizeof(M);
|
||||
}
|
||||
|
||||
__host__
|
||||
void this_will_not_compile_if_detector_is_marked_device_only()
|
||||
{
|
||||
float f = Wrapper<float, Functor<float>>().f<float>();
|
||||
inline float getRandomFloat(float min = 10, float max = 100) {
|
||||
std::uniform_real_distribution<float> gen(min, max);
|
||||
return gen(rng);
|
||||
}
|
||||
|
||||
template <typename T, typename B>
|
||||
void fillMatrix(T* a, int size) {
|
||||
for (int i = 0; i < size; i++) {
|
||||
T t;
|
||||
t.x = getRandomFloat();
|
||||
if constexpr (count<T, B>() >= 2) t.y = getRandomFloat();
|
||||
if constexpr (count<T, B>() >= 3) t.z = getRandomFloat();
|
||||
if constexpr (count<T, B>() >= 4) t.w = getRandomFloat();
|
||||
|
||||
a[i] = t;
|
||||
}
|
||||
}
|
||||
|
||||
// Test operations
|
||||
template <typename T, typename B>
|
||||
__host__ __device__ void testOperations(T& a, T& b) {
|
||||
a.x += b.x;
|
||||
a.x++;
|
||||
b.x++;
|
||||
if constexpr (count<T, B>() >= 2) {
|
||||
a.y = b.x;
|
||||
a.x = b.y;
|
||||
}
|
||||
if constexpr (count<T, B>() >= 3) {
|
||||
if (a.x > 0) b.x /= a.x;
|
||||
a.x *= b.z;
|
||||
a.y--;
|
||||
}
|
||||
if constexpr (count<T, B>() >= 4) {
|
||||
b.w = a.x;
|
||||
a.w += (-b.y);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T, typename B>
|
||||
__global__ void testOperationsGPU(T* d_a, T* d_b, int size) {
|
||||
int id = threadIdx.x;
|
||||
if (id > size) return;
|
||||
T &a = d_a[id];
|
||||
T &b = d_b[id];
|
||||
|
||||
testOperations<T, B>(a, b);
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
void dcopy(T* a, T* b, int size) {
|
||||
for (int i = 0; i < size; i++) {
|
||||
a[i] = b[i];
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
bool isEqual(T* a, T* b, int size) {
|
||||
for (int i = 0; i < size; i++) {
|
||||
if (a[i] != b[i]) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
// Main function that tests type
|
||||
// T = what you want to test
|
||||
// D = pack of 1 i.e. float1 int1
|
||||
template <typename T, typename D>
|
||||
void testType(int msize) {
|
||||
T *fa, *fb, *fc, *h_fa, *h_fb;
|
||||
fa = new T[msize];
|
||||
fb = new T[msize];
|
||||
fc = new T[msize];
|
||||
h_fa = new T[msize];
|
||||
h_fb = new T[msize];
|
||||
|
||||
T *d_fa, *d_fb;
|
||||
|
||||
constexpr int c = count<T, D>();
|
||||
|
||||
if (c <= 0 || c >= 5) {
|
||||
failed("Invalid Size\n");
|
||||
}
|
||||
|
||||
fillMatrix<T, D>(fa, msize);
|
||||
dcopy(fb, fa, msize);
|
||||
dcopy(h_fa, fa, msize);
|
||||
dcopy(h_fb, fa, msize);
|
||||
for (int i = 0; i < msize; i++) testOperations<T, D>(h_fa[i], h_fb[i]);
|
||||
|
||||
hipMalloc(&d_fa, sizeof(T) * msize);
|
||||
hipMalloc(&d_fb, sizeof(T) * msize);
|
||||
|
||||
hipMemcpy(d_fa, fa, sizeof(T) * msize, hipMemcpyHostToDevice);
|
||||
hipMemcpy(d_fb, fb, sizeof(T) * msize, hipMemcpyHostToDevice);
|
||||
|
||||
auto kernel = testOperationsGPU<T, D>;
|
||||
hipLaunchKernelGGL(kernel, 1, msize, 0, 0, d_fa, d_fb, msize);
|
||||
|
||||
hipMemcpy(fc, d_fa, sizeof(T) * msize, hipMemcpyDeviceToHost);
|
||||
|
||||
bool pass = true;
|
||||
if (!isEqual<T>(h_fa, fc, msize)) {
|
||||
pass = false;
|
||||
}
|
||||
|
||||
delete[] fa;
|
||||
delete[] fb;
|
||||
delete[] fc;
|
||||
delete[] h_fa;
|
||||
delete[] h_fb;
|
||||
hipFree(d_fa);
|
||||
hipFree(d_fb);
|
||||
|
||||
if (!pass) {
|
||||
failed("Failed");
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
const int msize = 100;
|
||||
// double
|
||||
testType<double1, double1>(msize);
|
||||
testType<double2, double1>(msize);
|
||||
testType<double3, double1>(msize);
|
||||
testType<double4, double1>(msize);
|
||||
|
||||
// floats
|
||||
testType<float1, float1>(msize);
|
||||
testType<float2, float1>(msize);
|
||||
testType<float3, float1>(msize);
|
||||
testType<float4, float1>(msize);
|
||||
...
|
||||
passed();
|
||||
}
|
||||
```
|
||||
[1]: https://en.wikibooks.org/wiki/More_C%2B%2B_Idioms/Member_Detector
|
||||
For more details for the complete program, please refer to HIP test application at the link, https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/deviceLib/hip_floatnTM.cpp
|
||||
|
||||
@@ -56,10 +56,10 @@ and provides practical suggestions on how to port CUDA code and work through com
|
||||
- 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.
|
||||
- Use **[hipconvertinplace-perl.sh](https://github.com/ROCm-Developer-Tools/HIPIFY/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](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/bin/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.
|
||||
@@ -113,7 +113,7 @@ For each input file FILE, this script will:
|
||||
This is useful for testing improvements to the hipify toolset.
|
||||
|
||||
|
||||
The [hipconvertinplace-perl.sh](https://github.com/ROCm-Developer-Tools/HIP/blob/master/bin/hipconvertinplace-perl.sh) script will perform inplace conversion for all code files in the specified directory.
|
||||
The [hipconvertinplace-perl.sh](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/bin/hipconvertinplace-perl.sh) script will perform inplace conversion for all code files in the specified directory.
|
||||
This can be quite handy when dealing with an existing CUDA code base since the script preserves the existing directory structure
|
||||
and filenames - and includes work. After converting in-place, you can review the code to add additional parameters to
|
||||
directory names.
|
||||
|
||||
@@ -121,7 +121,8 @@ hipRTC APIs accept HIP source files in character string format as input paramete
|
||||
|
||||
For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://github.com/RadeonOpenCompute/ROCm).
|
||||
|
||||
The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism.
|
||||
The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/main/docs/markdown/hip_rtc.md).
|
||||
|
||||
|
||||
## Device-Side Malloc
|
||||
|
||||
|
||||
@@ -12,22 +12,22 @@ detailed overview, see the help text available with `roc-obj --help`.
|
||||
### Examples:
|
||||
|
||||
#### Extract all ROCm code objects from a list of executables
|
||||
roc-obj executable...
|
||||
roc-obj <executable>...
|
||||
|
||||
#### Extract all ROCm code objects from a list of executables, and disassemble them
|
||||
roc-obj --disassemble executable...
|
||||
roc-obj --disassemble <executable>...
|
||||
# or
|
||||
roc-obj -d executable...
|
||||
roc-obj -d <executable>...
|
||||
|
||||
#### Extract all ROCm code objects from a list of executables into dir/
|
||||
roc-obj --outdir dir/ executable...
|
||||
roc-obj --outdir dir/ <executable>...
|
||||
# or
|
||||
roc-obj -o dir/ executable...
|
||||
roc-obj -o dir/ <executable>...
|
||||
|
||||
#### Extract only ROCm code objects matching regex over Target ID
|
||||
roc-obj --grep gfx9 executable...
|
||||
roc-obj --target-id gfx9 <executable>...
|
||||
# or
|
||||
roc-obj -g gfx9 executable...
|
||||
roc-obj -t gfx9 <executable>...
|
||||
|
||||
## Low-Level Tooling
|
||||
|
||||
@@ -48,19 +48,19 @@ detailed overview, see the help text available with `roc-obj --help`.
|
||||
|
||||
### List available ROCm Code Objects: rocm-obj-ls
|
||||
|
||||
Use this tool to list available ROCm code objects. Code objects are listed using URI syntax.
|
||||
Use this tool to list available ROCm code objects. Code objects are listed by bundle number, entry ID, and URI syntax.
|
||||
|
||||
Usage: roc-obj-ls [-v|h] executable...
|
||||
List the URIs of the code objects embedded in the specfied host executables.
|
||||
-v Verbose output (includes Entry ID)
|
||||
-v Verbose output. Adds column headers for more human readable format
|
||||
-h Show this help message
|
||||
|
||||
|
||||
### Extract ROCm Code Objects: rocm-obj-extract
|
||||
### Extract ROCm Code Objects: roc-obj-extract
|
||||
|
||||
Extracts available ROCm code objects from specified URI.
|
||||
|
||||
Usage: rocm-obj-extract [-o|v|h] URI...
|
||||
Usage: roc-obj-extract [-o|v|h] URI...
|
||||
- URIs can be read from STDIN, one per line.
|
||||
- From the URIs specified, extracts code objects into files named: <executable_name>-[pid<number>]-offset<number>-size<number>.co
|
||||
|
||||
@@ -74,14 +74,11 @@ detailed overview, see the help text available with `roc-obj --help`.
|
||||
|
||||
### Examples:
|
||||
|
||||
#### Dump all code objects to current directory:
|
||||
roc-obj-ls <exe> | roc-obj-extract
|
||||
|
||||
#### Dump the ISA for gfx906:
|
||||
roc-obj-ls -v <exe> | awk '/gfx906/{print $2}' | roc-obj-extract -o - | llvm-objdump -d - > <exe>.gfx906.isa
|
||||
roc-obj-ls -v <exe> | awk '/gfx906/{print $3}' | roc-obj-extract -o - | llvm-objdump -d - > <exe>.gfx906.isa
|
||||
|
||||
#### Check the e_flags of the gfx908 code object:
|
||||
roc-obj-ls -v <exe> | awk '/gfx908/{print $2}' | roc-obj-extract -o - | llvm-readelf -h - | grep Flags
|
||||
roc-obj-ls -v <exe> | awk '/gfx908/{print $3}' | roc-obj-extract -o - | llvm-readelf -h - | grep Flags
|
||||
|
||||
#### Disassemble the fourth code object:
|
||||
roc-obj-ls <exe> | sed -n 4p | roc-obj-extract -o - | llvm-objdump -d -
|
||||
@@ -90,6 +87,6 @@ detailed overview, see the help text available with `roc-obj --help`.
|
||||
for uri in $(roc-obj-ls <exe>); do printf "%d: %s\n" "$(roc-obj-extract -o - "$uri" | wc -c)" "$uri"; done | sort -n
|
||||
|
||||
#### Compare disassembly of gfx803 and gfx900 code objects:
|
||||
dis() { roc-obj-ls -v <exe> | grep "$1" | awk '{print $2}' | roc-obj-extract -o - | llvm-objdump -d -; }
|
||||
dis() { roc-obj-ls -v <exe> | grep "$1" | awk '{print $3}' | roc-obj-extract -o - | llvm-objdump -d -; }
|
||||
diff <(dis gfx803) <(dis gfx900)
|
||||
|
||||
|
||||
@@ -87,6 +87,7 @@ endif()
|
||||
find_path(HSA_HEADER hsa/hsa.h
|
||||
PATHS
|
||||
"${_IMPORT_PREFIX}/../include"
|
||||
"${ROCM_PATH}/include"
|
||||
/opt/rocm/include
|
||||
)
|
||||
|
||||
@@ -94,6 +95,15 @@ if (HSA_HEADER-NOTFOUND)
|
||||
message (FATAL_ERROR "HSA header not found! ROCM_PATH environment not set")
|
||||
endif()
|
||||
|
||||
file(GLOB HIP_CLANGRT_LIB_SEARCH_PATHS "${CMAKE_HIP_COMPILER}/../lib/clang/*/lib/*")
|
||||
find_library(CLANGRT_BUILTINS
|
||||
NAMES
|
||||
clang_rt.builtins
|
||||
clang_rt.builtins-x86_64
|
||||
PATHS
|
||||
${HIP_CLANGRT_LIB_SEARCH_PATHS}
|
||||
${HIP_CLANG_INCLUDE_PATH}/../lib/linux)
|
||||
|
||||
set_target_properties(hip-lang::device PROPERTIES
|
||||
INTERFACE_INCLUDE_DIRECTORIES "$<$<COMPILE_LANGUAGE:HIP>:${_IMPORT_PREFIX}/../include;${HIP_CLANG_INCLUDE_PATH}>"
|
||||
INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "$<$<COMPILE_LANGUAGE:HIP>:${_IMPORT_PREFIX}/../include;${HIP_CLANG_INCLUDE_PATH}>"
|
||||
@@ -123,9 +133,13 @@ set_property(TARGET hip-lang::device APPEND PROPERTY
|
||||
)
|
||||
|
||||
# Add support for __fp16 and _Float16, explicitly link with compiler-rt
|
||||
set_property(TARGET hip-lang::device APPEND PROPERTY
|
||||
INTERFACE_LINK_LIBRARIES "$<$<LINK_LANGUAGE:HIP>:${HIP_CLANG_INCLUDE_PATH}/lib/linux/libclang_rt.builtins-x86_64.a>"
|
||||
)
|
||||
if(CLANGRT_BUILTINS-NOTFOUND)
|
||||
message(FATAL_ERROR "clangrt builtins lib not found")
|
||||
else()
|
||||
set_property(TARGET hip-lang::device APPEND PROPERTY
|
||||
INTERFACE_LINK_LIBRARIES "$<$<LINK_LANGUAGE:HIP>:${CLANGRT_BUILTINS}>"
|
||||
)
|
||||
endif()
|
||||
|
||||
# Approved by CMake to use this name. This is used so that HIP can
|
||||
# change the name of the target and not require any modifications in CMake
|
||||
|
||||
+118
-96
@@ -3931,6 +3931,66 @@ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3
|
||||
* @{
|
||||
* This section describes the texture management functions of HIP runtime API.
|
||||
*/
|
||||
hipError_t hipBindTextureToMipmappedArray(
|
||||
const textureReference* tex,
|
||||
hipMipmappedArray_const_t mipmappedArray,
|
||||
const hipChannelFormatDesc* desc);
|
||||
hipError_t hipGetTextureReference(
|
||||
const textureReference** texref,
|
||||
const void* symbol);
|
||||
hipError_t hipCreateTextureObject(
|
||||
hipTextureObject_t* pTexObject,
|
||||
const hipResourceDesc* pResDesc,
|
||||
const hipTextureDesc* pTexDesc,
|
||||
const struct hipResourceViewDesc* pResViewDesc);
|
||||
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
|
||||
hipError_t hipGetChannelDesc(
|
||||
hipChannelFormatDesc* desc,
|
||||
hipArray_const_t array);
|
||||
hipError_t hipGetTextureObjectResourceDesc(
|
||||
hipResourceDesc* pResDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
hipError_t hipGetTextureObjectResourceViewDesc(
|
||||
struct hipResourceViewDesc* pResViewDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
hipError_t hipGetTextureObjectTextureDesc(
|
||||
hipTextureDesc* pTexDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
hipError_t hipTexRefSetAddressMode(
|
||||
textureReference* texRef,
|
||||
int dim,
|
||||
enum hipTextureAddressMode am);
|
||||
hipError_t hipTexRefSetArray(
|
||||
textureReference* tex,
|
||||
hipArray_const_t array,
|
||||
unsigned int flags);
|
||||
hipError_t hipTexRefSetFilterMode(
|
||||
textureReference* texRef,
|
||||
enum hipTextureFilterMode fm);
|
||||
hipError_t hipTexRefSetFlags(
|
||||
textureReference* texRef,
|
||||
unsigned int Flags);
|
||||
hipError_t hipTexRefSetFormat(
|
||||
textureReference* texRef,
|
||||
hipArray_Format fmt,
|
||||
int NumPackedComponents);
|
||||
hipError_t hipTexObjectCreate(
|
||||
hipTextureObject_t* pTexObject,
|
||||
const HIP_RESOURCE_DESC* pResDesc,
|
||||
const HIP_TEXTURE_DESC* pTexDesc,
|
||||
const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
|
||||
hipError_t hipTexObjectDestroy(
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetResourceDesc(
|
||||
HIP_RESOURCE_DESC* pResDesc,
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetResourceViewDesc(
|
||||
HIP_RESOURCE_VIEW_DESC* pResViewDesc,
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetTextureDesc(
|
||||
HIP_TEXTURE_DESC* pTexDesc,
|
||||
hipTextureObject_t texObject);
|
||||
|
||||
/**
|
||||
*
|
||||
* @addtogroup TexturD Texture Management [Deprecated]
|
||||
@@ -3965,35 +4025,6 @@ hipError_t hipGetTextureAlignmentOffset(
|
||||
const textureReference* texref);
|
||||
DEPRECATED(DEPRECATED_MSG)
|
||||
hipError_t hipUnbindTexture(const textureReference* tex);
|
||||
// doxygen end deprecated texture management
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
hipError_t hipBindTextureToMipmappedArray(
|
||||
const textureReference* tex,
|
||||
hipMipmappedArray_const_t mipmappedArray,
|
||||
const hipChannelFormatDesc* desc);
|
||||
hipError_t hipGetTextureReference(
|
||||
const textureReference** texref,
|
||||
const void* symbol);
|
||||
hipError_t hipCreateTextureObject(
|
||||
hipTextureObject_t* pTexObject,
|
||||
const hipResourceDesc* pResDesc,
|
||||
const hipTextureDesc* pTexDesc,
|
||||
const struct hipResourceViewDesc* pResViewDesc);
|
||||
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
|
||||
hipError_t hipGetChannelDesc(
|
||||
hipChannelFormatDesc* desc,
|
||||
hipArray_const_t array);
|
||||
hipError_t hipGetTextureObjectResourceDesc(
|
||||
hipResourceDesc* pResDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
hipError_t hipGetTextureObjectResourceViewDesc(
|
||||
struct hipResourceViewDesc* pResViewDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
hipError_t hipGetTextureObjectTextureDesc(
|
||||
hipTextureDesc* pTexDesc,
|
||||
hipTextureObject_t textureObject);
|
||||
DEPRECATED(DEPRECATED_MSG)
|
||||
hipError_t hipTexRefGetAddress(
|
||||
hipDeviceptr_t* dev_ptr,
|
||||
@@ -4049,49 +4080,23 @@ hipError_t hipTexRefSetAddress2D(
|
||||
const HIP_ARRAY_DESCRIPTOR* desc,
|
||||
hipDeviceptr_t dptr,
|
||||
size_t Pitch);
|
||||
hipError_t hipTexRefSetAddressMode(
|
||||
textureReference* texRef,
|
||||
int dim,
|
||||
enum hipTextureAddressMode am);
|
||||
hipError_t hipTexRefSetArray(
|
||||
textureReference* tex,
|
||||
hipArray_const_t array,
|
||||
unsigned int flags);
|
||||
hipError_t hipTexRefSetFilterMode(
|
||||
textureReference* texRef,
|
||||
enum hipTextureFilterMode fm);
|
||||
hipError_t hipTexRefSetFlags(
|
||||
textureReference* texRef,
|
||||
unsigned int Flags);
|
||||
hipError_t hipTexRefSetFormat(
|
||||
textureReference* texRef,
|
||||
hipArray_Format fmt,
|
||||
int NumPackedComponents);
|
||||
DEPRECATED(DEPRECATED_MSG)
|
||||
hipError_t hipTexRefSetMaxAnisotropy(
|
||||
textureReference* texRef,
|
||||
unsigned int maxAniso);
|
||||
hipError_t hipTexObjectCreate(
|
||||
hipTextureObject_t* pTexObject,
|
||||
const HIP_RESOURCE_DESC* pResDesc,
|
||||
const HIP_TEXTURE_DESC* pTexDesc,
|
||||
const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
|
||||
hipError_t hipTexObjectDestroy(
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetResourceDesc(
|
||||
HIP_RESOURCE_DESC* pResDesc,
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetResourceViewDesc(
|
||||
HIP_RESOURCE_VIEW_DESC* pResViewDesc,
|
||||
hipTextureObject_t texObject);
|
||||
hipError_t hipTexObjectGetTextureDesc(
|
||||
HIP_TEXTURE_DESC* pTexDesc,
|
||||
hipTextureObject_t texObject);
|
||||
// doxygen end Texture management
|
||||
// doxygen end deprecated texture management
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
// The following are not supported.
|
||||
/**
|
||||
*
|
||||
* @addtogroup TextureU Texture Management [Not supported]
|
||||
* @{
|
||||
* @ingroup Texture
|
||||
* This section describes the texture management functions currently unsupported in HIP runtime.
|
||||
*/
|
||||
DEPRECATED(DEPRECATED_MSG)
|
||||
hipError_t hipTexRefSetBorderColor(
|
||||
textureReference* texRef,
|
||||
@@ -4120,6 +4125,15 @@ hipError_t hipMipmappedArrayGetLevel(
|
||||
hipArray_t* pLevelArray,
|
||||
hipMipmappedArray_t hMipMappedArray,
|
||||
unsigned int level);
|
||||
// doxygen end Texture management unsupported
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
// doxygen end Texture management
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
@@ -4136,7 +4150,10 @@ hipError_t hipMipmappedArrayGetLevel(
|
||||
*/
|
||||
|
||||
/**
|
||||
* Callback/Activity API
|
||||
*
|
||||
* @defgroup Callback Callback Activity APIs
|
||||
* @{
|
||||
* This section describes the callback/Activity of HIP runtime API.
|
||||
*/
|
||||
hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg);
|
||||
hipError_t hipRemoveApiCallback(uint32_t id);
|
||||
@@ -4147,6 +4164,10 @@ const char* hipKernelNameRef(const hipFunction_t f);
|
||||
const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream);
|
||||
int hipGetStreamDeviceId(hipStream_t stream);
|
||||
|
||||
// doxygen end Callback
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
@@ -5160,6 +5181,38 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
|
||||
* @}
|
||||
*/
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup GL Interop
|
||||
* @{
|
||||
* This section describes Stream Memory Wait and Write functions of HIP runtime API.
|
||||
*/
|
||||
typedef unsigned int GLuint;
|
||||
|
||||
// Queries devices associated with GL Context.
|
||||
hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
|
||||
unsigned int hipDeviceCount, hipGLDeviceList deviceList);
|
||||
// Registers a GL Buffer for interop and returns corresponding graphics resource.
|
||||
hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer,
|
||||
unsigned int flags);
|
||||
// Maps a graphics resource for hip access.
|
||||
hipError_t hipGraphicsMapResources(int count, hipGraphicsResource_t* resources,
|
||||
hipStream_t stream __dparm(0) );
|
||||
// Gets device accessible address of a graphics resource.
|
||||
hipError_t hipGraphicsResourceGetMappedPointer(void** devPtr, size_t* size,
|
||||
hipGraphicsResource_t resource);
|
||||
// Unmaps a graphics resource for hip access.
|
||||
hipError_t hipGraphicsUnmapResources(int count, hipGraphicsResource_t* resources,
|
||||
hipStream_t stream __dparm(0));
|
||||
// Unregisters a graphics resource.
|
||||
hipError_t hipGraphicsUnregisterResource(hipGraphicsResource_t resource);
|
||||
// doxygen end GL Interop
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
} /* extern "c" */
|
||||
#endif
|
||||
@@ -5331,37 +5384,6 @@ static inline hipError_t hipUnbindTexture(
|
||||
|
||||
#endif // __cplusplus
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup GL Interop
|
||||
* @{
|
||||
* This section describes Stream Memory Wait and Write functions of HIP runtime API.
|
||||
*/
|
||||
typedef unsigned int GLuint;
|
||||
|
||||
// Queries devices associated with GL Context.
|
||||
hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
|
||||
unsigned int hipDeviceCount, hipGLDeviceList deviceList);
|
||||
// Registers a GL Buffer for interop and returns corresponding graphics resource.
|
||||
hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer,
|
||||
unsigned int flags);
|
||||
// Maps a graphics resource for hip access.
|
||||
hipError_t hipGraphicsMapResources(int count, hipGraphicsResource_t* resources,
|
||||
hipStream_t stream __dparm(0) );
|
||||
// Gets device accessible address of a graphics resource.
|
||||
hipError_t hipGraphicsResourceGetMappedPointer(void** devPtr, size_t* size,
|
||||
hipGraphicsResource_t resource);
|
||||
// Unmaps a graphics resource for hip access.
|
||||
hipError_t hipGraphicsUnmapResources(int count, hipGraphicsResource_t* resources,
|
||||
hipStream_t stream __dparm(0));
|
||||
// Unregisters a graphics resource.
|
||||
hipError_t hipGraphicsUnregisterResource(hipGraphicsResource_t resource);
|
||||
// doxygen end GL Interop
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
#ifdef __GNUC__
|
||||
#pragma GCC visibility pop
|
||||
#endif
|
||||
|
||||
@@ -26,12 +26,16 @@ if(NOT DEFINED __HIP_ENABLE_PCH)
|
||||
set(__HIP_ENABLE_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
|
||||
endif()
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
if(${__HIP_ENABLE_PCH})
|
||||
add_definitions(-D__HIP_ENABLE_PCH)
|
||||
endif()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -20,7 +20,8 @@
|
||||
|
||||
#Dependencies : [MYHIP]/bin must be in user's path.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(module_api)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(modile_api_global)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -24,11 +24,15 @@ project(square)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# create square.cpp
|
||||
execute_process(COMMAND sh -c "/opt/rocm/hip/bin/hipify-perl ../square.cu > ../square.cpp")
|
||||
execute_process(COMMAND sh -c "${ROCM_PATH}/hip/bin/hipify-perl ../square.cu > ../square.cpp")
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -19,9 +19,9 @@ $ export HIP_PATH=[MYHIP]
|
||||
```
|
||||
$ cd ~/hip/samples/0_Intro/square
|
||||
$ make
|
||||
/home/user/hip/bin/hipify-perl square.cu > square.cpp
|
||||
/home/user/hip/bin/hipcc square.cpp -o square.out
|
||||
/home/user/hip/bin/hipcc -use-staticlib square.cpp -o square.out.static
|
||||
/opt/rocm/hip/bin/hipify-perl square.cu > square.cpp
|
||||
/opt/rocm/hip/bin/hipcc square.cpp -o square.out
|
||||
/opt/rocm/hip/bin/hipcc -use-staticlib square.cpp -o square.out.static
|
||||
```
|
||||
- Execute file
|
||||
```
|
||||
|
||||
@@ -22,8 +22,12 @@ project(hipBusBandwidth)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(hipCommander)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(hipDispatchLatency)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(hipInfo)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -75,7 +75,7 @@ void printDeviceProp(int deviceId) {
|
||||
<< endl;
|
||||
cout << setw(w1) << "device#" << deviceId << endl;
|
||||
|
||||
hipDeviceProp_t props;
|
||||
hipDeviceProp_t props = {0};
|
||||
HIPCHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
|
||||
cout << setw(w1) << "Name: " << props.name << endl;
|
||||
@@ -90,12 +90,8 @@ void printDeviceProp(int deviceId) {
|
||||
cout << setw(w1) << "memoryClockRate: " << (float)props.memoryClockRate / 1000.0 << " Mhz"
|
||||
<< endl;
|
||||
cout << setw(w1) << "memoryBusWidth: " << props.memoryBusWidth << endl;
|
||||
cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0
|
||||
<< " Mhz" << endl;
|
||||
cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2)
|
||||
<< bytesToGB(props.totalGlobalMem) << " GB" << endl;
|
||||
cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2)
|
||||
<< bytesToKB(props.maxSharedMemoryPerMultiProcessor) << " KB" << endl;
|
||||
cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl;
|
||||
cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB"
|
||||
<< endl;
|
||||
@@ -116,6 +112,21 @@ void printDeviceProp(int deviceId) {
|
||||
cout << setw(w1) << "concurrentKernels: " << props.concurrentKernels << endl;
|
||||
cout << setw(w1) << "cooperativeLaunch: " << props.cooperativeLaunch << endl;
|
||||
cout << setw(w1) << "cooperativeMultiDeviceLaunch: " << props.cooperativeMultiDeviceLaunch << endl;
|
||||
cout << setw(w1) << "isIntegrated: " << props.integrated << endl;
|
||||
cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl;
|
||||
cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl;
|
||||
cout << setw(w1) << "maxTexture2D.height: " << props.maxTexture2D[1] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.width: " << props.maxTexture3D[0] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.height: " << props.maxTexture3D[1] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.depth: " << props.maxTexture3D[2] << endl;
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
cout << setw(w1) << "isLargeBar: " << props.isLargeBar << endl;
|
||||
cout << setw(w1) << "asicRevision: " << props.asicRevision << endl;
|
||||
cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2)
|
||||
<< bytesToKB(props.maxSharedMemoryPerMultiProcessor) << " KB" << endl;
|
||||
cout << setw(w1) << "clockInstructionRate: " << (float)props.clockInstructionRate / 1000.0
|
||||
<< " Mhz" << endl;
|
||||
cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl;
|
||||
cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch
|
||||
<< endl;
|
||||
@@ -136,16 +147,7 @@ void printDeviceProp(int deviceId) {
|
||||
cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl;
|
||||
cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl;
|
||||
cout << setw(w1) << "gcnArchName: " << props.gcnArchName << endl;
|
||||
cout << setw(w1) << "isIntegrated: " << props.integrated << endl;
|
||||
cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl;
|
||||
cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl;
|
||||
cout << setw(w1) << "maxTexture2D.height: " << props.maxTexture2D[1] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.width: " << props.maxTexture3D[0] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.height: " << props.maxTexture3D[1] << endl;
|
||||
cout << setw(w1) << "maxTexture3D.depth: " << props.maxTexture3D[2] << endl;
|
||||
cout << setw(w1) << "isLargeBar: " << props.isLargeBar << endl;
|
||||
cout << setw(w1) << "asicRevision: " << props.asicRevision << endl;
|
||||
|
||||
#endif
|
||||
int deviceCnt;
|
||||
hipGetDeviceCount(&deviceCnt);
|
||||
cout << setw(w1) << "peers: ";
|
||||
|
||||
@@ -22,8 +22,12 @@ project(MatrixTranspose)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(inline_asm)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(texture2dDrv)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -19,10 +19,13 @@
|
||||
# THE SOFTWARE.
|
||||
|
||||
cmake_minimum_required(VERSION 2.8.3)
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
if(NOT DEFINED HIP_PATH)
|
||||
if(NOT DEFINED ENV{HIP_PATH})
|
||||
set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to which HIP has been installed")
|
||||
set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH "Path to which HIP has been installed")
|
||||
else()
|
||||
set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed")
|
||||
endif()
|
||||
@@ -36,7 +39,7 @@ find_package(HIP QUIET)
|
||||
if(HIP_FOUND)
|
||||
message(STATUS "Found HIP: " ${HIP_VERSION})
|
||||
else()
|
||||
message(FATAL_ERROR "Could not find HIP. Ensure that HIP is either installed in /opt/rocm/hip or the variable HIP_PATH is set to point to the right location.")
|
||||
message(FATAL_ERROR "Could not find HIP. Ensure that HIP is either installed in ${ROCM_PATH}/hip or the variable HIP_PATH is set to point to the right location.")
|
||||
endif()
|
||||
|
||||
set(MY_SOURCE_FILES MatrixTranspose.cpp)
|
||||
@@ -50,7 +53,7 @@ set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FO
|
||||
hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${HIP_PATH} /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${HIP_PATH} ${ROCM_PATH})
|
||||
find_package(hip QUIET)
|
||||
if(TARGET hip::host)
|
||||
message(STATUS "Found hip::host at ${hip_DIR}")
|
||||
|
||||
@@ -22,8 +22,12 @@ project(occupancy)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(gpuarch)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -74,7 +74,8 @@ hipcc hipMain1.cpp -L. -lHipOptLibrary -o test_emit_static_hipcc_linker.out
|
||||
```
|
||||
Using g++:
|
||||
```
|
||||
g++ hipMain1.cpp -L. -lHipOptLibrary -L/opt/rocm/hip/lib -lamdhip64 -o test_emit_static_host_linker.out
|
||||
ROCM_PATH is the path where ROCM is installed. default path is /opt/rocm.
|
||||
g++ hipMain1.cpp -L. -lHipOptLibrary -L<ROCM_PATH>/hip/lib -lamdhip64 -o test_emit_static_host_linker.out
|
||||
```
|
||||
|
||||
## Static libraries with device functions
|
||||
|
||||
@@ -2,8 +2,12 @@ project(static_lib)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip REQUIRED)
|
||||
@@ -23,7 +27,7 @@ add_library(HipDevice STATIC ${CPP_SOURCES})
|
||||
|
||||
target_compile_options(HipDevice PRIVATE -fgpu-rdc)
|
||||
target_link_libraries(HipDevice PRIVATE -fgpu-rdc)
|
||||
target_include_directories(HipDevice PRIVATE /opt/rocm/hsa/include)
|
||||
target_include_directories(HipDevice PRIVATE ${ROCM_PATH}/hsa/include)
|
||||
|
||||
# Create test executable that uses libHipDevice.a
|
||||
set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain2.cpp)
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -2,8 +2,12 @@ project(static_lib)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip REQUIRED)
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -1,3 +1,4 @@
|
||||
ROCM_PATH is the path where ROCM is installed. default path is /opt/rocm.
|
||||
# Compile to assembly and create an executable from modified asm
|
||||
|
||||
This sample shows how to generate the assembly code for a simple HIP source application, then re-compiling it and generating a valid HIP executable.
|
||||
@@ -7,8 +8,8 @@ This sample uses a previous HIP application sample, please see [0_Intro/square](
|
||||
## Compiling the HIP source into assembly
|
||||
Using HIP flags `-c -S` will help generate the host x86_64 and the device AMDGCN assembly code when paired with `--cuda-host-only` and `--cuda-device-only` respectively. In this sample we use these commands:
|
||||
```
|
||||
/opt/rocm/hip/bin/hipcc -c -S --cuda-host-only -target x86_64-linux-gnu -o square_host.s square.cpp
|
||||
/opt/rocm/hip/bin/hipcc -c -S --cuda-device-only --offload-arch=gfx900 --offload-arch=gfx906 square.cpp
|
||||
<ROCM_PATH>/hip/bin/hipcc -c -S --cuda-host-only -target x86_64-linux-gnu -o square_host.s square.cpp
|
||||
<ROCM_PATH>/hip/bin/hipcc -c -S --cuda-device-only --offload-arch=gfx900 --offload-arch=gfx906 square.cpp
|
||||
```
|
||||
|
||||
The device assembly will be output into two separate files:
|
||||
@@ -22,22 +23,22 @@ You may modify `--offload-arch` flag to build other archs and choose to enable o
|
||||
## Compiling the assembly into a valid HIP executable
|
||||
If valid, the modified host and device assembly may be compiled into a HIP executable. The host assembly can be compiled into an object using this command:
|
||||
```
|
||||
/opt/rocm/hip/bin/hipcc -c square_host.s -o square_host.o
|
||||
<ROCM_PATH>/hip/bin/hipcc -c square_host.s -o square_host.o
|
||||
```
|
||||
|
||||
However, the device assembly code will require a few extra steps. The device assemblies needs to be compiled into device objects, then offload-bundled into a HIP fat binary using the clang-offload-bundler, then llvm-mc embeds the binary inside of a host object using the MC directives provided in `hip_obj_gen.mcin`. The output is a host object with an embedded device object. Here are the steps for device side compilation into an object:
|
||||
```
|
||||
/opt/rocm/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx900 square-hip-amdgcn-amd-amdhsa-gfx900.s -o square-hip-amdgcn-amd-amdhsa-gfx900.o
|
||||
/opt/rocm/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 square-hip-amdgcn-amd-amdhsa-gfx906.s -o square-hip-amdgcn-amd-amdhsa-gfx906.o
|
||||
/opt/rocm/llvm/bin/clang-offload-bundler -type=o -bundle-align=4096 -targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa-gfx900,hip-amdgcn-amd-amdhsa-gfx906 -inputs=/dev/null,square-hip-amdgcn-amd-amdhsa-gfx900.o,square-hip-amdgcn-amd-amdhsa-gfx906.o -outputs=offload_bundle.hipfb
|
||||
/opt/rocm/llvm/bin/llvm-mc -triple x86_64-unknown-linux-gnu hip_obj_gen.mcin -o square_device.o --filetype=obj
|
||||
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx900 square-hip-amdgcn-amd-amdhsa-gfx900.s -o square-hip-amdgcn-amd-amdhsa-gfx900.o
|
||||
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 square-hip-amdgcn-amd-amdhsa-gfx906.s -o square-hip-amdgcn-amd-amdhsa-gfx906.o
|
||||
<ROCM_PATH>/llvm/bin/clang-offload-bundler -type=o -bundle-align=4096 -targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa-gfx900,hip-amdgcn-amd-amdhsa-gfx906 -inputs=/dev/null,square-hip-amdgcn-amd-amdhsa-gfx900.o,square-hip-amdgcn-amd-amdhsa-gfx906.o -outputs=offload_bundle.hipfb
|
||||
<ROCM_PATH>/llvm/bin/llvm-mc -triple x86_64-unknown-linux-gnu hip_obj_gen.mcin -o square_device.o --filetype=obj
|
||||
```
|
||||
|
||||
**Note:** Using option `-bundle-align=4096` only works on ROCm 4.0 and newer compilers. Also, the architecture must match the same arch as when compiling to assembly.
|
||||
|
||||
Finally, using the system linker, hipcc, or clang, link the host and device objects into an executable:
|
||||
```
|
||||
/opt/rocm/hip/bin/hipcc square_host.o square_device.o -o square_asm.out
|
||||
<ROCM_PATH>/hip/bin/hipcc square_host.o square_device.o -o square_asm.out
|
||||
```
|
||||
If you haven't modified the GPU archs, this executable should run on both `gfx900` and `gfx906`.
|
||||
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -1,3 +1,4 @@
|
||||
ROCM_PATH is the path where ROCM is installed. default path is /opt/rocm.
|
||||
# Compile to LLVM IR and create an executable from modified IR
|
||||
|
||||
This sample shows how to generate the LLVM IR for a simple HIP source application, then re-compiling it and generating a valid HIP executable.
|
||||
@@ -7,8 +8,8 @@ This sample uses a previous HIP application sample, please see [0_Intro/square](
|
||||
## Compiling the HIP source into LLVM IR
|
||||
Using HIP flags `-c -emit-llvm` will help generate the host x86_64 and the device LLVM bitcode when paired with `--cuda-host-only` and `--cuda-device-only` respectively. In this sample we use these commands:
|
||||
```
|
||||
/opt/rocm/hip/bin/hipcc -c -emit-llvm --cuda-host-only -target x86_64-linux-gnu -o square_host.bc square.cpp
|
||||
/opt/rocm/hip/bin/hipcc -c -emit-llvm --cuda-device-only --offload-arch=gfx900 --offload-arch=gfx906 square.cpp
|
||||
<ROCM_PATH>/hip/bin/hipcc -c -emit-llvm --cuda-host-only -target x86_64-linux-gnu -o square_host.bc square.cpp
|
||||
<ROCM_PATH>/hip/bin/hipcc -c -emit-llvm --cuda-device-only --offload-arch=gfx900 --offload-arch=gfx906 square.cpp
|
||||
```
|
||||
The device LLVM IR bitcode will be output into two separate files:
|
||||
- square-hip-amdgcn-amd-amdhsa-gfx900.bc
|
||||
@@ -18,8 +19,8 @@ You may modify `--offload-arch` flag to build other archs and choose to enable o
|
||||
|
||||
To transform the LLVM bitcode into human readable LLVM IR, use these commands:
|
||||
```
|
||||
/opt/rocm/llvm/bin/llvm-dis square-hip-amdgcn-amd-amdhsa-gfx900.bc -o square-hip-amdgcn-amd-amdhsa-gfx900.ll
|
||||
/opt/rocm/llvm/bin/llvm-dis square-hip-amdgcn-amd-amdhsa-gfx906.bc -o square-hip-amdgcn-amd-amdhsa-gfx906.ll
|
||||
<ROCM_PATH>/llvm/bin/llvm-dis square-hip-amdgcn-amd-amdhsa-gfx900.bc -o square-hip-amdgcn-amd-amdhsa-gfx900.ll
|
||||
<ROCM_PATH>/llvm/bin/llvm-dis square-hip-amdgcn-amd-amdhsa-gfx906.bc -o square-hip-amdgcn-amd-amdhsa-gfx906.ll
|
||||
```
|
||||
|
||||
**Warning:** We cannot ensure any compiler besides the ROCm hipcc and clang will be compatible with this process. Also, there is no guarantee that the starting IR produced with `-x cl` will run with HIP runtime. Experimenting with other compilers or starting IR will be the responsibility of the developer.
|
||||
@@ -32,25 +33,25 @@ At this point, you may evaluate the LLVM IR and make modifications if you are fa
|
||||
## Compiling the LLVM IR into a valid HIP executable
|
||||
If valid, the modified host and device IR may be compiled into a HIP executable. First, the readable IR must be compiled back in LLVM bitcode. The host IR can be compiled into an object using this command:
|
||||
```
|
||||
/opt/rocm/llvm/bin/llvm-as square_host.ll -o square_host.bc
|
||||
/opt/rocm/hip/bin/hipcc -c square_host.bc -o square_host.o
|
||||
<ROCM_PATH>/llvm/bin/llvm-as square_host.ll -o square_host.bc
|
||||
<ROCM_PATH>/hip/bin/hipcc -c square_host.bc -o square_host.o
|
||||
```
|
||||
|
||||
However, the device IR will require a few extra steps. The device bitcodes needs to be compiled into device objects, then offload-bundled into a HIP fat binary using the clang-offload-bundler, then llvm-mc embeds the binary inside of a host object using the MC directives provided in `hip_obj_gen.mcin`. The output is a host object with an embedded device object. Here are the steps for device side compilation into an object:
|
||||
```
|
||||
/opt/rocm/hip/../llvm/bin/llvm-as square-hip-amdgcn-amd-amdhsa-gfx900.ll -o square-hip-amdgcn-amd-amdhsa-gfx900.bc
|
||||
/opt/rocm/hip/../llvm/bin/llvm-as square-hip-amdgcn-amd-amdhsa-gfx906.ll -o square-hip-amdgcn-amd-amdhsa-gfx906.bc
|
||||
/opt/rocm/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx900 square-hip-amdgcn-amd-amdhsa-gfx900.bc -o square-hip-amdgcn-amd-amdhsa-gfx900.o
|
||||
/opt/rocm/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 square-hip-amdgcn-amd-amdhsa-gfx906.bc -o square-hip-amdgcn-amd-amdhsa-gfx906.o
|
||||
/opt/rocm/hip/../llvm/bin/clang-offload-bundler -type=o -bundle-align=4096 -targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa-gfx900,hip-amdgcn-amd-amdhsa-gfx906 -inputs=/dev/null,square-hip-amdgcn-amd-amdhsa-gfx900.o,square-hip-amdgcn-amd-amdhsa-gfx906.o -outputs=offload_bundle.hipfb
|
||||
/opt/rocm/llvm/bin/llvm-mc hip_obj_gen.mcin -o square_device.o --filetype=obj
|
||||
<ROCM_PATH>/hip/../llvm/bin/llvm-as square-hip-amdgcn-amd-amdhsa-gfx900.ll -o square-hip-amdgcn-amd-amdhsa-gfx900.bc
|
||||
<ROCM_PATH>/hip/../llvm/bin/llvm-as square-hip-amdgcn-amd-amdhsa-gfx906.ll -o square-hip-amdgcn-amd-amdhsa-gfx906.bc
|
||||
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx900 square-hip-amdgcn-amd-amdhsa-gfx900.bc -o square-hip-amdgcn-amd-amdhsa-gfx900.o
|
||||
<ROCM_PATH>/hip/../llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 square-hip-amdgcn-amd-amdhsa-gfx906.bc -o square-hip-amdgcn-amd-amdhsa-gfx906.o
|
||||
<ROCM_PATH>/hip/../llvm/bin/clang-offload-bundler -type=o -bundle-align=4096 -targets=host-x86_64-unknown-linux,hip-amdgcn-amd-amdhsa-gfx900,hip-amdgcn-amd-amdhsa-gfx906 -inputs=/dev/null,square-hip-amdgcn-amd-amdhsa-gfx900.o,square-hip-amdgcn-amd-amdhsa-gfx906.o -outputs=offload_bundle.hipfb
|
||||
<ROCM_PATH>/llvm/bin/llvm-mc hip_obj_gen.mcin -o square_device.o --filetype=obj
|
||||
```
|
||||
|
||||
**Note:** Using option `-bundle-align=4096` only works on ROCm 4.0 and newer compilers. Also, the architecture must match the same arch as when compiling to LLVM IR.
|
||||
|
||||
Finally, using the system linker, hipcc, or clang, link the host and device objects into an executable:
|
||||
```
|
||||
/opt/rocm/hip/bin/hipcc square_host.o square_device.o -o square_ir.out
|
||||
<ROCM_PATH>/hip/bin/hipcc square_host.o square_device.o -o square_ir.out
|
||||
```
|
||||
If you haven't modified the GPU archs, this executable should run on both `gfx900` and `gfx906`.
|
||||
|
||||
|
||||
@@ -22,8 +22,12 @@ project(hipEvent)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(sharedMemory)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(shfl)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(2dshfl)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(dynamic_shared)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(stream)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(peer2peer)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
@@ -22,8 +22,12 @@ project(unroll)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
@@ -18,7 +18,8 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
|
||||
+1
-1
@@ -24,7 +24,7 @@ cmake ..
|
||||
|
||||
b. to build with static libs, run
|
||||
|
||||
cmake -DCMAKE_PREFIX_PATH="/opt/rocm/llvm/lib/cmake" ..
|
||||
cmake -DCMAKE_PREFIX_PATH="<ROCM_PATH>/llvm/lib/cmake" ..
|
||||
|
||||
Then run,
|
||||
|
||||
|
||||
@@ -3,8 +3,6 @@ set(TEST_SRC
|
||||
add.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(ABMAddKernels SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests ABMAddKernels)
|
||||
hip_add_exe_to_target(NAME ABMAddKernels
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
|
||||
@@ -1,3 +1,19 @@
|
||||
cmake_minimum_required(VERSION 3.16.8)
|
||||
project(hiptests)
|
||||
|
||||
# Check if platform and compiler are set
|
||||
if(HIP_PLATFORM STREQUAL "amd")
|
||||
if(HIP_COMPILER STREQUAL "nvcc")
|
||||
message(FATAL_ERROR "Unexpected HIP_COMPILER:${HIP_COMPILER} is set for HIP_PLATFOR:amd")
|
||||
endif()
|
||||
elseif(HIP_PLATFORM STREQUAL "nvidia")
|
||||
if(NOT DEFINED HIP_COMPILER OR NOT HIP_COMPILER STREQUAL "nvcc")
|
||||
message(FATAL_ERROR "Unexpected HIP_COMPILER: ${HIP_COMPILER} is set for HIP_PLATFORM:nvidia")
|
||||
endif()
|
||||
else()
|
||||
message(FATAL_ERROR "Unexpected HIP_PLATFORM: " ${HIP_PLATFORM})
|
||||
endif()
|
||||
|
||||
# Set HIP Path
|
||||
if(NOT DEFINED HIP_PATH)
|
||||
if(DEFINED ENV{HIP_PATH})
|
||||
@@ -8,7 +24,22 @@ if(NOT DEFINED HIP_PATH)
|
||||
endif()
|
||||
message(STATUS "HIP Path: ${HIP_PATH}")
|
||||
|
||||
set(CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc")
|
||||
if(UNIX)
|
||||
set(CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc")
|
||||
set(CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc")
|
||||
else()
|
||||
# using cmake_path as it handles path correctly.
|
||||
# Set both compilers else windows cmake complains of mismatch
|
||||
cmake_path(SET CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc.bat")
|
||||
cmake_path(SET CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc.bat")
|
||||
endif()
|
||||
|
||||
|
||||
if(NOT UNIX)
|
||||
# In linux this reruns the cmake and fails with incorrect vars.
|
||||
# so the project command is used only for windows
|
||||
project(build_tests)
|
||||
endif()
|
||||
|
||||
if(NOT DEFINED CATCH2_PATH)
|
||||
if(DEFINED ENV{CATCH2_PATH})
|
||||
@@ -102,10 +133,10 @@ add_subdirectory(unit)
|
||||
add_subdirectory(ABM)
|
||||
add_subdirectory(hipTestMain)
|
||||
add_subdirectory(stress)
|
||||
add_subdirectory(TypeQualifiers)
|
||||
|
||||
if(UNIX)
|
||||
add_subdirectory(multiproc)
|
||||
add_subdirectory(TypeQualifiers)
|
||||
endif()
|
||||
|
||||
cmake_policy(POP)
|
||||
|
||||
@@ -3,10 +3,6 @@ set(TEST_SRC
|
||||
hipManagedKeyword.cc
|
||||
)
|
||||
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(TypeQualifiers SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests TypeQualifiers)
|
||||
|
||||
hip_add_exe_to_target(NAME TypeQualifiers
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
|
||||
+42
-1
@@ -173,7 +173,7 @@ function(catch_discover_tests TARGET)
|
||||
"if(EXISTS \"${ctest_tests_file}\")\n"
|
||||
" include(\"${ctest_tests_file}\")\n"
|
||||
"else()\n"
|
||||
" add_test(${TARGET}_NOT_BUILT-${args_hash} ${TARGET}_NOT_BUILT-${args_hash})\n"
|
||||
" message(WARNING \"Test ${TARGET} not built yet.\")\n"
|
||||
"endif()\n"
|
||||
)
|
||||
|
||||
@@ -204,3 +204,44 @@ set(_CATCH_DISCOVER_TESTS_SCRIPT
|
||||
${CMAKE_CURRENT_LIST_DIR}/CatchAddTests.cmake
|
||||
CACHE INTERNAL "Catch2 full path to CatchAddTests.cmake helper file"
|
||||
)
|
||||
|
||||
###############################################################################
|
||||
# function to be called by all tests
|
||||
function(hip_add_exe_to_target)
|
||||
set(options)
|
||||
set(args NAME TEST_TARGET_NAME PLATFORM COMPILE_OPTIONS)
|
||||
set(list_args TEST_SRC LINKER_LIBS PROPERTY)
|
||||
cmake_parse_arguments(
|
||||
PARSE_ARGV 0
|
||||
"" # variable prefix
|
||||
"${options}"
|
||||
"${args}"
|
||||
"${list_args}"
|
||||
)
|
||||
# Create shared lib of all tests
|
||||
add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object>)
|
||||
catch_discover_tests(${_NAME} PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
if(UNIX)
|
||||
set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs)
|
||||
endif()
|
||||
if(DEFINED _LINKER_LIBS)
|
||||
target_link_libraries(${_NAME} ${_LINKER_LIBS})
|
||||
endif()
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(${_TEST_TARGET_NAME} ${_NAME})
|
||||
|
||||
if (DEFINED _PROPERTY)
|
||||
set_property(TARGET ${_NAME} PROPERTY ${_PROPERTY})
|
||||
endif()
|
||||
|
||||
if (DEFINED _COMPILE_OPTIONS)
|
||||
target_compile_options(${_NAME} PUBLIC ${_COMPILE_OPTIONS})
|
||||
endif()
|
||||
|
||||
foreach(arg IN LISTS _UNPARSED_ARGUMENTS)
|
||||
message(WARNING "Unparsed arguments: ${arg}")
|
||||
endforeach()
|
||||
endfunction()
|
||||
|
||||
|
||||
|
||||
@@ -22,83 +22,9 @@ if(CMAKE_BUILD_TYPE MATCHES "^Debug$")
|
||||
add_definitions(-DHT_LOG_ENABLE)
|
||||
endif()
|
||||
|
||||
add_executable(UnitTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
add_library(Main_Object EXCLUDE_FROM_ALL OBJECT main.cc hip_test_context.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET UnitTests PROPERTY CXX_STANDARD 17)
|
||||
set_property(TARGET Main_Object PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(UnitTests PUBLIC -std=c++17)
|
||||
target_compile_options(Main_Object PUBLIC -std=c++17)
|
||||
endif()
|
||||
|
||||
target_link_libraries(UnitTests PRIVATE UnitDeviceTests
|
||||
MemoryTest
|
||||
StreamTest
|
||||
EventTest
|
||||
OccupancyTest
|
||||
DeviceTest
|
||||
RTC
|
||||
printfTests
|
||||
TextureTest
|
||||
stdc++fs)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "nvidia")
|
||||
target_link_libraries(UnitTests PRIVATE nvrtc)
|
||||
endif()
|
||||
|
||||
catch_discover_tests(UnitTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
|
||||
# ABM exe
|
||||
add_executable(ABMTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET ABMTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(ABMTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
|
||||
target_link_libraries(ABMTests PRIVATE ABMAddKernels
|
||||
stdc++fs)
|
||||
|
||||
catch_discover_tests(ABMTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
|
||||
add_dependencies(build_tests UnitTests ABMTests)
|
||||
|
||||
# Add Multiproc tests as seperate binary
|
||||
if(UNIX)
|
||||
add_executable(MultiProcTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET MultiProcTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(MultiProcTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
|
||||
target_link_libraries(MultiProcTests PRIVATE MultiProc
|
||||
stdc++fs)
|
||||
catch_discover_tests(MultiProcTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
add_dependencies(build_tests MultiProcTests)
|
||||
endif()
|
||||
|
||||
add_executable(StressTest EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
add_custom_target(build_stress_test)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET StressTest PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(StressTest PUBLIC -std=c++17)
|
||||
endif()
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
target_link_libraries(StressTest PRIVATE printf stream)
|
||||
endif()
|
||||
target_link_libraries(StressTest PRIVATE memory stdc++fs)
|
||||
add_dependencies(build_stress_test StressTest)
|
||||
add_custom_target(stress_test COMMAND StressTest)
|
||||
|
||||
# Space Specifiers/Qualifiers exe
|
||||
add_executable(TypeQualifierTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set_property(TARGET TypeQualifierTests PROPERTY CXX_STANDARD 17)
|
||||
else()
|
||||
target_compile_options(TypeQualifierTests PUBLIC -std=c++17)
|
||||
endif()
|
||||
target_link_libraries(TypeQualifierTests PRIVATE TypeQualifiers stdc++fs)
|
||||
|
||||
catch_discover_tests(TypeQualifierTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
|
||||
add_dependencies(build_tests TypeQualifierTests)
|
||||
|
||||
@@ -2,18 +2,8 @@
|
||||
#include <picojson.h>
|
||||
#include <fstream>
|
||||
#include <sstream>
|
||||
|
||||
#if __has_include(<filesystem>)
|
||||
#include <filesystem>
|
||||
namespace fs = std::filesystem;
|
||||
#elif __has_include(<experimental/filesystem>)
|
||||
#include <experimental/filesystem>
|
||||
namespace fs = std::experimental::filesystem;
|
||||
#else
|
||||
#error "gg filesystem"
|
||||
#endif
|
||||
|
||||
#include <regex>
|
||||
#include "hip_test_filesystem.hh"
|
||||
|
||||
void TestContext::detectOS() {
|
||||
#if (HT_WIN == 1)
|
||||
@@ -49,7 +39,7 @@ void TestContext::fillConfig() {
|
||||
if (config_path.has_parent_path() && config_path.has_filename()) {
|
||||
config_.json_file = config_str;
|
||||
} else if (config_path.has_parent_path()) {
|
||||
config_.json_file = config_path / def_config_json;
|
||||
config_.json_file = config_path.string() + def_config_json;
|
||||
} else {
|
||||
config_.json_file = exe_path + def_config_json;
|
||||
}
|
||||
|
||||
@@ -23,6 +23,7 @@ THE SOFTWARE.
|
||||
#pragma once
|
||||
#include "hip_test_context.hh"
|
||||
#include <catch.hpp>
|
||||
#include <stdlib.h>
|
||||
|
||||
#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
|
||||
|
||||
@@ -104,4 +105,14 @@ static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlo
|
||||
|
||||
return blocks;
|
||||
}
|
||||
|
||||
static inline int RAND_R(unsigned* rand_seed)
|
||||
{
|
||||
#if defined(_WIN32) || defined(_WIN64)
|
||||
srand(*rand_seed);
|
||||
return rand();
|
||||
#else
|
||||
return rand_r(rand_seed);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,89 @@
|
||||
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
// We haven't checked which filesystem to include yet
|
||||
#ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL
|
||||
// Check for feature test macro for <filesystem>
|
||||
#if defined(__cpp_lib_filesystem)
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0
|
||||
// Check for feature test macro for <experimental/filesystem>
|
||||
#elif defined(__cpp_lib_experimental_filesystem)
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1
|
||||
// We can't check if headers exist...
|
||||
// Let's assume experimental to be safe
|
||||
#elif !defined(__has_include)
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1
|
||||
// Check if the header "<filesystem>" exists
|
||||
#elif __has_include(<filesystem>)
|
||||
// If we're compiling on Visual Studio and are not compiling with C++17,
|
||||
// we need to use experimental
|
||||
#ifdef _MSC_VER
|
||||
// Check and include header that defines "_HAS_CXX17"
|
||||
#if __has_include(<yvals_core.h>)
|
||||
#include <yvals_core.h>
|
||||
|
||||
// Check for enabled C++17 support
|
||||
#if defined(_HAS_CXX17) && _HAS_CXX17
|
||||
// We're using C++17, so let's use the normal version
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
// If the marco isn't defined yet, that means any of the other
|
||||
// VS specific checks failed, so we need to use experimental
|
||||
#ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1
|
||||
#endif
|
||||
|
||||
// Not on Visual Studio. Let's use the normal version
|
||||
#else // #ifdef _MSC_VER
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0
|
||||
#endif
|
||||
|
||||
// Check if the header "<filesystem>" exists
|
||||
#elif __has_include(<experimental/filesystem>)
|
||||
#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1
|
||||
|
||||
// Fail if neither header is available with a nice error message
|
||||
#else
|
||||
#error Could not find system header "<filesystem>" ||
|
||||
"<experimental/filesystem>"
|
||||
#endif
|
||||
|
||||
// We priously determined that we need the exprimental version
|
||||
#if INCLUDE_STD_FILESYSTEM_EXPERIMENTAL
|
||||
// Include it
|
||||
#define _SILENCE_EXPERIMENTAL_FILESYSTEM_DEPRECATION_WARNING 1;
|
||||
#include <experimental/filesystem>
|
||||
// We need the alias from std::experimental::filesystem to std::filesystem
|
||||
namespace fs = std::experimental::filesystem;
|
||||
// We have a decent compiler and can use the normal version
|
||||
#else
|
||||
// Include it
|
||||
#include <filesystem>
|
||||
namespace fs = std::filesystem;
|
||||
#endif
|
||||
|
||||
#endif // #ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL
|
||||
@@ -24,7 +24,10 @@ THE SOFTWARE.
|
||||
#include "hip_test_common.hh"
|
||||
|
||||
#ifdef __linux__
|
||||
#include <sys/sysinfo.h>
|
||||
#include <sys/sysinfo.h>
|
||||
#else
|
||||
#include <windows.h>
|
||||
#include <sysinfoapi.h>
|
||||
#endif
|
||||
|
||||
namespace HipTest {
|
||||
|
||||
@@ -29,16 +29,7 @@ THE SOFTWARE.
|
||||
#include <random>
|
||||
#include <fstream>
|
||||
#include <streambuf>
|
||||
|
||||
#if __has_include(<filesystem>)
|
||||
#include <filesystem>
|
||||
namespace fs = std::filesystem;
|
||||
#elif __has_include(<experimental/filesystem>)
|
||||
#include <experimental/filesystem>
|
||||
namespace fs = std::experimental::filesystem;
|
||||
#else
|
||||
#error "gg filesystem"
|
||||
#endif
|
||||
#include "hip_test_filesystem.hh"
|
||||
|
||||
namespace hip {
|
||||
class SpawnProc {
|
||||
|
||||
@@ -11,14 +11,13 @@ set(LINUX_TEST_SRC
|
||||
hipIpcMemAccessTest.cc
|
||||
hipHostMallocTestsMproc.cc
|
||||
hipMallocConcurrencyMproc.cc
|
||||
hipMemCoherencyTstMProc.cc
|
||||
)
|
||||
|
||||
if(UNIX)
|
||||
# Create shared lib of all tests
|
||||
add_library(MultiProc SHARED EXCLUDE_FROM_ALL ${LINUX_TEST_SRC})
|
||||
|
||||
target_link_libraries(MultiProc ${CMAKE_DL_LIBS})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests MultiProc)
|
||||
# the last argument linker libraries is required for this test but optional to the function
|
||||
hip_add_exe_to_target(NAME MultiProc
|
||||
TEST_SRC ${LINUX_TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
LINKER_LIBS ${CMAKE_DL_LIBS})
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,809 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario 3: The test validates if fine grain
|
||||
behavior is observed or not with memory allocated using malloc()
|
||||
Scenario 4: The test validates if coarse grain memory
|
||||
behavior is observed or not with memory allocated using malloc()
|
||||
Scenario 5: The test validates if fine memory
|
||||
behavior is observed or not with memory allocated using mmap()
|
||||
Scenario 6: The test validates if coarse grain memory
|
||||
behavior is observed or not with memory allocated using mmap()
|
||||
Scenario:7 Test Case Description: The following test checks if the memory is
|
||||
accessible when HIP_HOST_COHERENT is set to 0
|
||||
Scenario:8 Test Case Description: The following test checks if the memory
|
||||
exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <unistd.h>
|
||||
#include <sys/mman.h>
|
||||
#include <sys/wait.h>
|
||||
#include <chrono>
|
||||
|
||||
__global__ void CoherentTst(int *ptr, int PeakClk) {
|
||||
// Incrementing the value by 1
|
||||
int64_t GpuFrq = (PeakClk * 1000);
|
||||
int64_t StrtTck = clock64();
|
||||
atomicAdd(ptr, 1);
|
||||
// The following while loop checks the value in ptr for around 3-4 seconds
|
||||
while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
|
||||
if (*ptr == 3) {
|
||||
atomicAdd(ptr, 1);
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void SquareKrnl(int *ptr) {
|
||||
// ptr value squared here
|
||||
*ptr = (*ptr) * (*ptr);
|
||||
}
|
||||
|
||||
|
||||
|
||||
// The variable below will work as signal to decide pass/fail
|
||||
static bool YES_COHERENT = false;
|
||||
|
||||
// The function tests the coherency of allocated memory
|
||||
static void TstCoherency(int *Ptr, bool HmmMem) {
|
||||
int *Dptr = nullptr, peak_clk;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
// storing value 1 in the memory created above
|
||||
*Ptr = 1;
|
||||
// Getting gpu frequency
|
||||
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
|
||||
if (!HmmMem) {
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr), Ptr,
|
||||
0));
|
||||
CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk);
|
||||
} else {
|
||||
CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk);
|
||||
}
|
||||
// looping until the value is 2 for 3 seconds
|
||||
std::chrono::steady_clock::time_point start =
|
||||
std::chrono::steady_clock::now();
|
||||
while (std::chrono::duration_cast<std::chrono::seconds>(
|
||||
std::chrono::steady_clock::now() - start).count() < 3) {
|
||||
if (*Ptr == 2) {
|
||||
*Ptr += 1;
|
||||
break;
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 4) {
|
||||
YES_COHERENT = true;
|
||||
}
|
||||
}
|
||||
|
||||
/* Test case description: The following test validates if fine grain
|
||||
behavior is observed or not with memory allocated using malloc()*/
|
||||
// The following test is failing on Nvidia platform hence disabled it for now
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_malloc_CoherentTst") {
|
||||
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// The following code block is used to check for gfx906/8 so as to skip if
|
||||
// any of the gpus available
|
||||
int fd1[2]; // Used to store two ends of first pipe
|
||||
pid_t p;
|
||||
if (pipe(fd1) == -1) {
|
||||
fprintf(stderr, "Pipe Failed");
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
|
||||
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
|
||||
int GpuId[2] = {0, 0};
|
||||
p = fork();
|
||||
|
||||
if (p < 0) {
|
||||
fprintf(stderr, "fork Failed");
|
||||
REQUIRE(false);
|
||||
} else if (p > 0) { // parent process
|
||||
close(fd1[1]); // Close writing end of first pipe
|
||||
// Wait for child to send a string
|
||||
wait(NULL);
|
||||
// Read string from child and close reading end.
|
||||
read(fd1[0], GpuId, 2 * sizeof(int));
|
||||
close(fd1[0]);
|
||||
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
|
||||
WARN("This test is not applicable on MI60 & MI100."
|
||||
"Skipping the test!!");
|
||||
exit(0);
|
||||
}
|
||||
} else { // child process
|
||||
close(fd1[0]); // Close read end of first pipe
|
||||
hipDeviceProp_t prop;
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, 0));
|
||||
char *p = NULL;
|
||||
p = strstr(prop.gcnArchName, "gfx906");
|
||||
if (p) {
|
||||
WARN("gfx906 gpu found on this system!!");
|
||||
GpuId[0] = 1;
|
||||
}
|
||||
p = strstr(prop.gcnArchName, "gfx908");
|
||||
if (p) {
|
||||
WARN("gfx908 gpu found on this system!!");
|
||||
GpuId[1] = 1;
|
||||
}
|
||||
// Write concatenated string and close writing end
|
||||
write(fd1[1], GpuId, 2 * sizeof(int));
|
||||
close(fd1[1]);
|
||||
exit(0);
|
||||
}
|
||||
|
||||
// Test Case execution begins from here
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int managed = 0;
|
||||
HIPCHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
if (managed == 1) {
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
bool HmmMem = true;
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipMallocManaged() memory
|
||||
Ptr = reinterpret_cast<int*>(malloc(SIZE));
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
free(Ptr);
|
||||
if (YES_COHERENT) {
|
||||
// exit() with code 10 which indicates pass
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
/* Test case description: The following test validates if coarse grain memory
|
||||
behavior is observed or not with memory allocated using malloc()*/
|
||||
// The following test is failing on Nvidia platform hence disabling it for now
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
|
||||
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// The following code block is used to check for gfx906/8 so as to skip if
|
||||
// any of the gpus available
|
||||
int fd1[2]; // Used to store two ends of first pipe
|
||||
pid_t p;
|
||||
if (pipe(fd1) == -1) {
|
||||
fprintf(stderr, "Pipe Failed");
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
|
||||
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
|
||||
int GpuId[2] = {0, 0};
|
||||
p = fork();
|
||||
|
||||
if (p < 0) {
|
||||
fprintf(stderr, "fork Failed");
|
||||
REQUIRE(false);
|
||||
} else if (p > 0) { // parent process
|
||||
close(fd1[1]); // Close writing end of first pipe
|
||||
// Wait for child to send a string
|
||||
wait(NULL);
|
||||
// Read string from child and close reading end.
|
||||
read(fd1[0], GpuId, 2 * sizeof(int));
|
||||
close(fd1[0]);
|
||||
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
|
||||
WARN("This test is not applicable on MI60 & MI100."
|
||||
"Skipping the test!!");
|
||||
exit(0);
|
||||
}
|
||||
} else { // child process
|
||||
close(fd1[0]); // Close read end of first pipe
|
||||
hipDeviceProp_t prop;
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, 0));
|
||||
char *p = NULL;
|
||||
p = strstr(prop.gcnArchName, "gfx906");
|
||||
if (p) {
|
||||
WARN("gfx906 gpu found on this system!!");
|
||||
GpuId[0] = 1;
|
||||
}
|
||||
p = strstr(prop.gcnArchName, "gfx908");
|
||||
if (p) {
|
||||
WARN("gfx908 gpu found on this system!!");
|
||||
GpuId[1] = 1;
|
||||
}
|
||||
// Write concatenated string and close writing end
|
||||
write(fd1[1], GpuId, 2 * sizeof(int));
|
||||
close(fd1[1]);
|
||||
exit(0);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int managed = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
if (managed == 1) {
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipMallocManaged() memory
|
||||
Ptr = reinterpret_cast<int*>(malloc(SIZE));
|
||||
*Ptr = 4;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 16) {
|
||||
// exit() with code 10 which indicates pass
|
||||
free(Ptr);
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
free(Ptr);
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test case description: The following test validates if fine memory
|
||||
behavior is observed or not with memory allocated using mmap()*/
|
||||
// The following test is failing on Nvidia platform hence disabling it for now
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_mmap_CoherentTst") {
|
||||
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// The following code block is used to check for gfx906/8 so as to skip if
|
||||
// any of the gpus available
|
||||
int fd1[2]; // Used to store two ends of first pipe
|
||||
pid_t p;
|
||||
if (pipe(fd1) == -1) {
|
||||
fprintf(stderr, "Pipe Failed");
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
|
||||
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
|
||||
int GpuId[2] = {0, 0};
|
||||
p = fork();
|
||||
|
||||
if (p < 0) {
|
||||
fprintf(stderr, "fork Failed");
|
||||
REQUIRE(false);
|
||||
} else if (p > 0) { // parent process
|
||||
close(fd1[1]); // Close writing end of first pipe
|
||||
// Wait for child to send a string
|
||||
wait(NULL);
|
||||
// Read string from child and close reading end.
|
||||
read(fd1[0], GpuId, 2 * sizeof(int));
|
||||
close(fd1[0]);
|
||||
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
|
||||
WARN("This test is not applicable on MI60 & MI100."
|
||||
"Skipping the test!!");
|
||||
exit(0);
|
||||
}
|
||||
} else { // child process
|
||||
close(fd1[0]); // Close read end of first pipe
|
||||
hipDeviceProp_t prop;
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, 0));
|
||||
char *p = NULL;
|
||||
p = strstr(prop.gcnArchName, "gfx906");
|
||||
if (p) {
|
||||
WARN("gfx906 gpu found on this system!!");
|
||||
GpuId[0] = 1;
|
||||
}
|
||||
p = strstr(prop.gcnArchName, "gfx908");
|
||||
if (p) {
|
||||
WARN("gfx908 gpu found on this system!!");
|
||||
GpuId[1] = 1;
|
||||
}
|
||||
// Write concatenated string and close writing end
|
||||
write(fd1[1], GpuId, 2 * sizeof(int));
|
||||
close(fd1[1]);
|
||||
exit(0);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int managed = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
if (managed == 1) {
|
||||
bool HmmMem = true;
|
||||
int *Ptr = reinterpret_cast<int*>(mmap(NULL, sizeof(int),
|
||||
PROT_READ | PROT_WRITE,
|
||||
MAP_PRIVATE | MAP_ANONYMOUS, 0, 0));
|
||||
if (Ptr == MAP_FAILED) {
|
||||
WARN("Mapping Failed\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// Initializing the value with 1
|
||||
*Ptr = 1;
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
int err = munmap(Ptr, sizeof(int));
|
||||
if (err != 0) {
|
||||
WARN("munmap failed\n");
|
||||
}
|
||||
if (YES_COHERENT) {
|
||||
exit(10);
|
||||
} else {
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test case description: The following test validates if coarse grain memory
|
||||
behavior is observed or not with memory allocated using mmap()*/
|
||||
// The following test is failing on Nvidia platform hence disabling it for now
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
|
||||
if ((setenv("HSA_XNACK", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
// The following code block is used to check for gfx906/8 so as to skip if
|
||||
// any of the gpus available
|
||||
int fd1[2]; // Used to store two ends of first pipe
|
||||
pid_t p;
|
||||
if (pipe(fd1) == -1) {
|
||||
fprintf(stderr, "Pipe Failed");
|
||||
REQUIRE(false);
|
||||
}
|
||||
|
||||
/* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
|
||||
GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
|
||||
int GpuId[2] = {0, 0};
|
||||
p = fork();
|
||||
|
||||
if (p < 0) {
|
||||
fprintf(stderr, "fork Failed");
|
||||
REQUIRE(false);
|
||||
} else if (p > 0) { // parent process
|
||||
close(fd1[1]); // Close writing end of first pipe
|
||||
// Wait for child to send a string
|
||||
wait(NULL);
|
||||
// Read string from child and close reading end.
|
||||
read(fd1[0], GpuId, 2 * sizeof(int));
|
||||
close(fd1[0]);
|
||||
if ((GpuId[0] == 1) || (GpuId[0] == 1)) {
|
||||
WARN("This test is not applicable on MI60 & MI100."
|
||||
"Skipping the test!!");
|
||||
exit(0);
|
||||
}
|
||||
} else { // child process
|
||||
close(fd1[0]); // Close read end of first pipe
|
||||
hipDeviceProp_t prop;
|
||||
HIPCHECK(hipGetDeviceProperties(&prop, 0));
|
||||
char *p = NULL;
|
||||
p = strstr(prop.gcnArchName, "gfx906");
|
||||
if (p) {
|
||||
WARN("gfx906 gpu found on this system!!");
|
||||
GpuId[0] = 1;
|
||||
}
|
||||
p = strstr(prop.gcnArchName, "gfx908");
|
||||
if (p) {
|
||||
WARN("gfx908 gpu found on this system!!");
|
||||
GpuId[1] = 1;
|
||||
}
|
||||
// Write concatenated string and close writing end
|
||||
write(fd1[1], GpuId, 2 * sizeof(int));
|
||||
close(fd1[1]);
|
||||
exit(0);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int managed = 0;
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
if (managed == 1) {
|
||||
int SIZE = sizeof(int);
|
||||
int *Ptr = reinterpret_cast<int*>(mmap(NULL, SIZE,
|
||||
PROT_READ | PROT_WRITE,
|
||||
MAP_PRIVATE | MAP_ANONYMOUS, 0, 0));
|
||||
if (Ptr == MAP_FAILED) {
|
||||
WARN("Mapping Failed\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
|
||||
// Initializing the value with 9
|
||||
*Ptr = 9;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
bool IfTstPassed = false;
|
||||
if (*Ptr == 81) {
|
||||
IfTstPassed = true;
|
||||
}
|
||||
int err = munmap(Ptr, SIZE);
|
||||
if (err != 0) {
|
||||
WARN("munmap failed\n");
|
||||
}
|
||||
if (IfTstPassed) {
|
||||
exit(10);
|
||||
} else {
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test Case Description: The following test checks if the memory is
|
||||
accessible when HIP_HOST_COHERENT is set to 0*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
|
||||
*Ptr = 4;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&PtrD), Ptr, 0));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(PtrD);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 16) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test Case Description: The following test checks if the memory is
|
||||
accessible when HIP_HOST_COHERENT is set to 0*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
|
||||
*Ptr = 4;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&PtrD), Ptr, 0));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(PtrD);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 16) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test Case Description: The following test checks if the memory is
|
||||
accessible when HIP_HOST_COHERENT is set to 0*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
|
||||
*Ptr = 4;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&PtrD), Ptr, 0));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(PtrD);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 16) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test Case Description: The following test checks if the memory is
|
||||
accessible when HIP_HOST_COHERENT is set to 0*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) {
|
||||
int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNonCoherent));
|
||||
*Ptr = 4;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&PtrD), Ptr, 0));
|
||||
SquareKrnl<<<1, 1, 0, strm>>>(PtrD);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
if (*Ptr == 16) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else {
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
/* Test Case Description: The following test checks if the memory exhibits
|
||||
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) { // child process
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
bool HmmMem = false;
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
|
||||
*Ptr = 4;
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
if (YES_COHERENT) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else { // parent process
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
/* Test Case Description: The following test checks if the memory exhibits
|
||||
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) { // child process
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
bool HmmMem = false;
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
|
||||
*Ptr = 1;
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
if (YES_COHERENT) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else { // parent process
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Test Case Description: The following test checks if the memory exhibits
|
||||
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) { // child process
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
bool HmmMem = false;
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
|
||||
*Ptr = 4;
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
if (YES_COHERENT) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else { // parent process
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
/* Test Case Description: The following test checks if the memory exhibits
|
||||
fine grain behavior when HIP_HOST_COHERENT is set to 1*/
|
||||
// The following test is AMD specific test hence skipping for Nvidia
|
||||
#if HT_AMD
|
||||
TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
|
||||
if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
|
||||
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
|
||||
REQUIRE(false);
|
||||
}
|
||||
int stat = 0;
|
||||
if (fork() == 0) { // child process
|
||||
int *Ptr = nullptr, SIZE = sizeof(int);
|
||||
bool HmmMem = false;
|
||||
YES_COHERENT = false;
|
||||
// Allocating hipHostMalloc() memory
|
||||
HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
|
||||
*Ptr = 1;
|
||||
TstCoherency(Ptr, HmmMem);
|
||||
if (YES_COHERENT) {
|
||||
// exit() with code 10 which indicates pass
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(10);
|
||||
} else {
|
||||
// exit() with code 9 which indicates fail
|
||||
HIP_CHECK(hipHostFree(Ptr));
|
||||
exit(9);
|
||||
}
|
||||
} else { // parent process
|
||||
wait(&stat);
|
||||
int Result = WEXITSTATUS(stat);
|
||||
if (Result != 10) {
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
@@ -1,3 +1,5 @@
|
||||
add_custom_target(stress_test COMMAND "${CMAKE_CTEST_COMMAND}" -R "Stress_")
|
||||
|
||||
add_subdirectory(memory)
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
add_subdirectory(printf)
|
||||
|
||||
@@ -2,10 +2,9 @@
|
||||
set(TEST_SRC
|
||||
memcpy.cc
|
||||
hipMemcpyMThreadMSize.cc
|
||||
hipMallocManagedStress.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(memory SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_stress_test memory)
|
||||
hip_add_exe_to_target(NAME memory
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME stress_test)
|
||||
|
||||
@@ -0,0 +1,106 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// The following test case allocation, host access, device access of HMM
|
||||
// memory from size 1 to 10KB
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
|
||||
#define INCRMNT 10
|
||||
// Kernel function
|
||||
__global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr,
|
||||
size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
for (size_t i = index; i < n; i += stride) {
|
||||
Hmm[i] = Dptr[i] + INCRMNT;
|
||||
}
|
||||
}
|
||||
static bool IfTestPassed = true;
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
INFO("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
INFO("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
// The following test case allocation, host access, device access of HMM
|
||||
// memory from size 1 to 10KB
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_MultiSize") {
|
||||
IfTestPassed = true;
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
unsigned char *Hmm1 = nullptr, *Hmm2 = nullptr;
|
||||
int InitVal = 100, blockSize = 64, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
for (int i = 1; i < (1024*1024); ++i) {
|
||||
HIP_CHECK(hipMallocManaged(&Hmm1, i));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm2, i));
|
||||
for (int j = 0; j < i; ++j) {
|
||||
Hmm1[j] = InitVal;
|
||||
}
|
||||
dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypesC<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, i);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the results
|
||||
for (int k = 0; k < i; ++k) {
|
||||
if (Hmm2[k] != (InitVal + INCRMNT)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
DataMismatch = 0;
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
HIP_CHECK(hipFree(Hmm2));
|
||||
if (IfTestPassed == false) {
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -261,7 +261,7 @@ void Memcpy_And_verify(int NUM_ELM) {
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipMemcpy_multiDevice-AllAPIs", "",
|
||||
TEMPLATE_TEST_CASE("Stress_hipMemcpy_multiDevice-AllAPIs", "",
|
||||
char, int, size_t, long double) {
|
||||
auto diff_size = GENERATE(1, 5, 10, 100, 1024, 10*1024, 100*1024,
|
||||
1024*1024, 10*1024*1024, 100*1024*1024,
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
TEST_CASE("hipMalloc", "DifferentSizes") {
|
||||
TEST_CASE("Stress_hipMalloc", "DifferentSizes") {
|
||||
int* d_a = nullptr;
|
||||
SECTION("Size 10") {
|
||||
auto res = hipMalloc(&d_a, sizeof(10));
|
||||
|
||||
@@ -4,8 +4,6 @@ set(TEST_SRC
|
||||
Stress_printf_SimpleKernels.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(printf SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_stress_test printf)
|
||||
hip_add_exe_to_target(NAME printf
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME stress_test)
|
||||
|
||||
@@ -3,8 +3,6 @@ set(TEST_SRC
|
||||
Stress_hipStreamCreate.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(stream SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_stress_test stream)
|
||||
hip_add_exe_to_target(NAME stream
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME stress_test)
|
||||
|
||||
@@ -28,3 +28,4 @@ add_subdirectory(rtc)
|
||||
add_subdirectory(printf)
|
||||
add_subdirectory(printfExe)
|
||||
add_subdirectory(texture)
|
||||
add_subdirectory(graph)
|
||||
|
||||
@@ -17,9 +17,6 @@ set(TEST_SRC
|
||||
hipSetGetDevice.cc
|
||||
)
|
||||
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(DeviceTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests DeviceTest)
|
||||
hip_add_exe_to_target(NAME DeviceTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
|
||||
@@ -9,12 +9,18 @@ set(TEST_SRC
|
||||
brev.cc
|
||||
popc.cc
|
||||
ldg.cc
|
||||
syncthreadsand.cc
|
||||
syncthreadscount.cc
|
||||
syncthreadsor.cc
|
||||
threadfence_system.cc
|
||||
)
|
||||
|
||||
# skipped for windows compiler issue - Illegal instruction detected
|
||||
if(UNIX)
|
||||
set(TEST_SRC ${TEST_SRC}
|
||||
syncthreadsand.cc
|
||||
syncthreadscount.cc
|
||||
syncthreadsor.cc)
|
||||
endif()
|
||||
|
||||
|
||||
# AMD only tests
|
||||
set(AMD_TEST_SRC
|
||||
unsafeAtomicAdd.cc
|
||||
@@ -28,15 +34,12 @@ set(AMD_TEST_SRC
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
|
||||
set_source_files_properties(floatTM.cc PROPERTIES COMPILE_FLAGS -std=c++17)
|
||||
hip_add_exe_to_target(NAME UnitDeviceTests
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
elseif(HIP_PLATFORM MATCHES "nvidia")
|
||||
hip_add_exe_to_target(NAME UnitDeviceTests
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS --Wno-deprecated-declarations)
|
||||
endif()
|
||||
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(UnitDeviceTests SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
if(HIP_PLATFORM MATCHES "nvidia")
|
||||
target_compile_options(UnitDeviceTests PUBLIC --Wno-deprecated-declarations)
|
||||
endif()
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests UnitDeviceTests)
|
||||
|
||||
@@ -3,12 +3,15 @@ set(TEST_SRC
|
||||
Unit_hipEvent_Negative.cc
|
||||
Unit_hipEvent.cc
|
||||
Unit_hipEventElapsedTime.cc
|
||||
Unit_hipEventRecord.cc
|
||||
Unit_hipEventIpc.cc
|
||||
)
|
||||
|
||||
# Create shared lib of all tests
|
||||
add_library(EventTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
# skipped for windows due to duplicate symbols
|
||||
if(UNIX)
|
||||
set(TEST_SRC ${TEST_SRC}
|
||||
Unit_hipEventRecord.cc
|
||||
Unit_hipEventIpc.cc)
|
||||
endif()
|
||||
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests EventTest)
|
||||
hip_add_exe_to_target(NAME EventTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
|
||||
Executable → Regular
+13
-10
@@ -1,5 +1,4 @@
|
||||
#!/bin/bash
|
||||
# Copyright (c) 2017 - 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
# Copyright (c) 2021 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
|
||||
@@ -19,13 +18,17 @@
|
||||
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
# THE SOFTWARE.
|
||||
|
||||
#usage : hipexamine-perl.sh DIRNAME [hipify-perl options]
|
||||
# Common Tests - Test independent of all platforms
|
||||
set(TEST_SRC
|
||||
hipGraphAddEmptyNode.cc
|
||||
hipGraphAddDependencies.cc
|
||||
hipGraph.cc
|
||||
hipSimpleGraphWithKernel.cc
|
||||
hipGraphAddMemcpyNode.cc
|
||||
)
|
||||
|
||||
# Generate HIP stats (LOC, CUDA->API conversions, missing functionality) for all the code files
|
||||
# in the specified directory.
|
||||
# Create shared lib of all tests
|
||||
add_library(GraphsTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
|
||||
|
||||
|
||||
SCRIPT_DIR=`dirname $0`
|
||||
SEARCH_DIR=$1
|
||||
shift
|
||||
$SCRIPT_DIR/hipify-perl -no-output -print-stats "$@" `$SCRIPT_DIR/findcode.sh $SEARCH_DIR`
|
||||
# Add dependency on build_tests to build it on this custom target
|
||||
add_dependencies(build_tests GraphsTest)
|
||||
@@ -0,0 +1,346 @@
|
||||
/*
|
||||
Copyright (c) 2021 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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
1) Execution Without HIPGraphs : Regular procedure of using stream with async api calls.
|
||||
2) Manual HIPGraph : Manual procedure of adding nodes to graphs and mapping dependencies.
|
||||
3) HIPGraphs Using StreamCapture : Capturing sequence of operations in stream and launching
|
||||
graph with the nodes automatically added.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#define THREADS_PER_BLOCK 512
|
||||
#define GRAPH_LAUNCH_ITERATIONS 1000
|
||||
|
||||
static __global__ void reduce(float* d_in, double* d_out) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
int tid = threadIdx.x;
|
||||
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) {
|
||||
d_in[myId] += d_in[myId + s];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
if (tid == 0) {
|
||||
d_out[blockIdx.x] = d_in[myId];
|
||||
}
|
||||
}
|
||||
|
||||
static __global__ void reduceFinal(double* d_in, double* d_out) {
|
||||
int myId = threadIdx.x + blockDim.x * blockIdx.x;
|
||||
int tid = threadIdx.x;
|
||||
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
|
||||
if (tid < s) {
|
||||
d_in[myId] += d_in[myId + s];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
if (tid == 0) {
|
||||
*d_out = d_in[myId];
|
||||
}
|
||||
}
|
||||
|
||||
static void init_input(float* a, size_t size) {
|
||||
unsigned int seed = time(nullptr);
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
a[i] = (rand_r(&seed) & 0xFF) / static_cast<float>(RAND_MAX);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Regular procedure of using stream with async api calls
|
||||
*/
|
||||
static void hipWithoutGraphs(float* inputVec_h, float* inputVec_d,
|
||||
double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) {
|
||||
hipStream_t stream1, stream2, stream3;
|
||||
hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2;
|
||||
double result_h = 0.0;
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
HIP_CHECK(hipStreamCreate(&stream3));
|
||||
HIP_CHECK(hipEventCreate(&forkStreamEvent));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent1));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent2));
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) {
|
||||
HIP_CHECK(hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * inputSize,
|
||||
hipMemcpyDefault, stream1));
|
||||
HIP_CHECK(hipMemsetAsync(outputVec_d, 0, sizeof(double) * numOfBlocks,
|
||||
stream2));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent1, stream2));
|
||||
HIP_CHECK(hipMemsetAsync(result_d, 0, sizeof(double), stream3));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent2, stream3));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0));
|
||||
hipLaunchKernelGGL(reduce, dim3(inputSize / THREADS_PER_BLOCK, 1, 1),
|
||||
dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, inputVec_d,
|
||||
outputVec_d);
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0));
|
||||
hipLaunchKernelGGL(reduceFinal, dim3(1, 1, 1),
|
||||
dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1,
|
||||
outputVec_d, result_d);
|
||||
HIP_CHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double),
|
||||
hipMemcpyDefault, stream1));
|
||||
HIP_CHECK(hipStreamSynchronize(stream1));
|
||||
}
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
auto result = std::chrono::duration<double, std::milli>(stop - start);
|
||||
INFO("Time taken for hipWithoutGraphs : "
|
||||
<< std::chrono::duration_cast<std::chrono::milliseconds>(result).count()
|
||||
<< " millisecs ");
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
HIP_CHECK(hipStreamDestroy(stream3));
|
||||
double result_h_cpu = 0.0;
|
||||
for (size_t i = 0; i < inputSize; i++) {
|
||||
result_h_cpu += inputVec_h[i];
|
||||
}
|
||||
|
||||
REQUIRE(result_h_cpu == result_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Capturing sequence of operations in stream and launching graph
|
||||
* with the nodes automatically added.
|
||||
*/
|
||||
static void hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d,
|
||||
double* outputVec_d, double* result_d,
|
||||
size_t inputSize, size_t numOfBlocks) {
|
||||
hipStream_t stream1, stream2, stream3, streamForGraph;
|
||||
hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2;
|
||||
hipGraph_t graph;
|
||||
double result_h = 0.0;
|
||||
HIP_CHECK(hipStreamCreate(&stream1));
|
||||
HIP_CHECK(hipStreamCreate(&stream2));
|
||||
HIP_CHECK(hipStreamCreate(&stream3));
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
HIP_CHECK(hipEventCreate(&forkStreamEvent));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent1));
|
||||
HIP_CHECK(hipEventCreate(&memsetEvent2));
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
|
||||
HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream3, forkStreamEvent, 0));
|
||||
HIP_CHECK(hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * inputSize,
|
||||
hipMemcpyDefault, stream1));
|
||||
HIP_CHECK(hipMemsetAsync(outputVec_d, 0, sizeof(double) * numOfBlocks,
|
||||
stream2));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent1, stream2));
|
||||
HIP_CHECK(hipMemsetAsync(result_d, 0, sizeof(double), stream3));
|
||||
HIP_CHECK(hipEventRecord(memsetEvent2, stream3));
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0));
|
||||
hipLaunchKernelGGL(reduce, dim3(inputSize / THREADS_PER_BLOCK, 1, 1),
|
||||
dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1,
|
||||
inputVec_d, outputVec_d);
|
||||
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0));
|
||||
hipLaunchKernelGGL(reduceFinal, dim3(1, 1, 1), dim3(THREADS_PER_BLOCK, 1, 1),
|
||||
0, stream1, outputVec_d, result_d);
|
||||
HIP_CHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double),
|
||||
hipMemcpyDefault, stream1));
|
||||
HIP_CHECK(hipStreamEndCapture(stream1, &graph));
|
||||
hipGraphNode_t* nodes{nullptr};
|
||||
size_t numNodes = 0;
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
|
||||
INFO("Num of nodes in the graph created using stream capture API"
|
||||
<< numNodes);
|
||||
HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numNodes));
|
||||
INFO("Num of root nodes in the graph created using"
|
||||
" stream capture API" << numNodes);
|
||||
hipGraphExec_t graphExec;
|
||||
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
auto start1 = std::chrono::high_resolution_clock::now();
|
||||
for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) {
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
auto withInit =
|
||||
std::chrono::duration<double, std::milli>(stop - start);
|
||||
auto withoutInit =
|
||||
std::chrono::duration<double, std::milli>(stop - start1);
|
||||
INFO("Time taken for hipGraphsUsingStreamCapture with Init: "
|
||||
<< std::chrono::duration_cast<std::chrono::milliseconds>(withInit).count()
|
||||
<< " milliseconds without Init:"
|
||||
<< std::chrono::duration_cast<std::chrono::milliseconds>(withoutInit).count()
|
||||
<< " milliseconds ");
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(stream1));
|
||||
HIP_CHECK(hipStreamDestroy(stream2));
|
||||
HIP_CHECK(hipStreamDestroy(stream3));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
double result_h_cpu = 0.0;
|
||||
for (size_t i = 0; i < inputSize; i++) {
|
||||
result_h_cpu += inputVec_h[i];
|
||||
}
|
||||
|
||||
REQUIRE(result_h_cpu == result_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Manual procedure of adding nodes to graphs and mapping dependencies.
|
||||
*/
|
||||
static void hipGraphsManual(float* inputVec_h, float* inputVec_d,
|
||||
double* outputVec_d, double* result_d, size_t inputSize,
|
||||
size_t numOfBlocks) {
|
||||
hipStream_t streamForGraph;
|
||||
hipGraph_t graph;
|
||||
std::vector<hipGraphNode_t> nodeDependencies;
|
||||
hipGraphNode_t memcpyNode, kernelNode, memsetNode;
|
||||
double result_h = 0.0;
|
||||
HIP_CHECK(hipStreamCreate(&streamForGraph));
|
||||
auto start = std::chrono::high_resolution_clock::now();
|
||||
hipKernelNodeParams kernelNodeParams{};
|
||||
hipMemsetParams memsetParams{};
|
||||
memsetParams.dst = reinterpret_cast<void*>(outputVec_d);
|
||||
memsetParams.value = 0;
|
||||
memsetParams.pitch = 0;
|
||||
memsetParams.elementSize = sizeof(float);
|
||||
memsetParams.width = numOfBlocks * 2;
|
||||
memsetParams.height = 1;
|
||||
HIP_CHECK(hipGraphCreate(&graph, 0));
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, inputVec_d,
|
||||
inputVec_h, sizeof(float) * inputSize, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr,
|
||||
0, &memsetParams));
|
||||
nodeDependencies.push_back(memsetNode);
|
||||
nodeDependencies.push_back(memcpyNode);
|
||||
void* kernelArgs[4] = {reinterpret_cast<void*>(&inputVec_d),
|
||||
reinterpret_cast<void*>(&outputVec_d), &inputSize,
|
||||
&numOfBlocks};
|
||||
kernelNodeParams.func = reinterpret_cast<void*>(reduce);
|
||||
kernelNodeParams.gridDim = dim3(inputSize / THREADS_PER_BLOCK, 1, 1);
|
||||
kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
|
||||
kernelNodeParams.sharedMemBytes = 0;
|
||||
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
|
||||
kernelNodeParams.extra = nullptr;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(),
|
||||
nodeDependencies.size(), &kernelNodeParams));
|
||||
nodeDependencies.clear();
|
||||
nodeDependencies.push_back(kernelNode);
|
||||
memset(&memsetParams, 0, sizeof(memsetParams));
|
||||
memsetParams.dst = result_d;
|
||||
memsetParams.value = 0;
|
||||
memsetParams.elementSize = sizeof(float);
|
||||
memsetParams.width = 2;
|
||||
memsetParams.height = 1;
|
||||
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
|
||||
&memsetParams));
|
||||
nodeDependencies.push_back(memsetNode);
|
||||
memset(&kernelNodeParams, 0, sizeof(kernelNodeParams));
|
||||
kernelNodeParams.func = reinterpret_cast<void*>(reduceFinal);
|
||||
kernelNodeParams.gridDim = dim3(1, 1, 1);
|
||||
kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1);
|
||||
kernelNodeParams.sharedMemBytes = 0;
|
||||
void* kernelArgs2[3] = {reinterpret_cast<void*>(&outputVec_d),
|
||||
reinterpret_cast<void*>(&result_d), &numOfBlocks};
|
||||
kernelNodeParams.kernelParams = kernelArgs2;
|
||||
kernelNodeParams.extra = nullptr;
|
||||
HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(),
|
||||
nodeDependencies.size(), &kernelNodeParams));
|
||||
nodeDependencies.clear();
|
||||
nodeDependencies.push_back(kernelNode);
|
||||
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph,
|
||||
nodeDependencies.data(), nodeDependencies.size(), &result_h,
|
||||
result_d, sizeof(double), hipMemcpyDeviceToHost));
|
||||
nodeDependencies.clear();
|
||||
nodeDependencies.push_back(memcpyNode);
|
||||
hipGraphExec_t graphExec;
|
||||
hipGraphNode_t* nodes{nullptr};
|
||||
size_t numNodes{};
|
||||
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
|
||||
INFO("Num of nodes in the graph created using hipGraphs Manual"
|
||||
<< numNodes);
|
||||
HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numNodes));
|
||||
INFO("Num of root nodes in the graph created using"
|
||||
" hipGraphs Manual" << numNodes);
|
||||
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
|
||||
auto start1 = std::chrono::high_resolution_clock::now();
|
||||
for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) {
|
||||
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
|
||||
}
|
||||
HIP_CHECK(hipStreamSynchronize(streamForGraph));
|
||||
auto stop = std::chrono::high_resolution_clock::now();
|
||||
auto withInit =
|
||||
std::chrono::duration<double, std::milli>(stop - start);
|
||||
auto withoutInit =
|
||||
std::chrono::duration<double, std::milli>(stop - start1);
|
||||
|
||||
INFO("Time taken for hipGraphsManual with Init: "
|
||||
<< std::chrono::duration_cast<std::chrono::milliseconds>(withInit).count()
|
||||
<< " milliseconds without Init:"
|
||||
<< std::chrono::duration_cast<std::chrono::milliseconds>(withoutInit).count()
|
||||
<< " milliseconds ");
|
||||
|
||||
HIP_CHECK(hipGraphExecDestroy(graphExec));
|
||||
HIP_CHECK(hipGraphDestroy(graph));
|
||||
HIP_CHECK(hipStreamDestroy(streamForGraph));
|
||||
double result_h_cpu = 0.0;
|
||||
for (size_t i = 0; i < inputSize; i++) {
|
||||
result_h_cpu += inputVec_h[i];
|
||||
}
|
||||
|
||||
REQUIRE(result_h_cpu == result_h);
|
||||
}
|
||||
|
||||
/**
|
||||
* Tests basic functionality of hipGraph APIs by
|
||||
* Execution Without HIPGraphs, Manual HIPGraph, HIPGraphs Using StreamCapture.
|
||||
*/
|
||||
TEST_CASE("Unit_hipGraph_BasicFunctional") {
|
||||
constexpr size_t size = 1 << 12;
|
||||
constexpr size_t maxBlocks = 512;
|
||||
float *inputVec_d{nullptr}, *inputVec_h{nullptr};
|
||||
double *outputVec_d{nullptr}, *result_d{nullptr};
|
||||
|
||||
INFO("Elements : " << size << " ThreadsPerBlock : " << THREADS_PER_BLOCK);
|
||||
INFO("Graph Launch iterations = " << GRAPH_LAUNCH_ITERATIONS);
|
||||
|
||||
hipSetDevice(0);
|
||||
inputVec_h = reinterpret_cast<float*>(malloc(sizeof(float) * size));
|
||||
REQUIRE(inputVec_h != nullptr);
|
||||
HIP_CHECK(hipMalloc(&inputVec_d, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&outputVec_d, sizeof(double) * maxBlocks));
|
||||
HIP_CHECK(hipMalloc(&result_d, sizeof(double)));
|
||||
init_input(inputVec_h, size);
|
||||
|
||||
SECTION("Execution Without HIPGraphs") {
|
||||
hipWithoutGraphs(inputVec_h, inputVec_d, outputVec_d,
|
||||
result_d, size, maxBlocks);
|
||||
}
|
||||
|
||||
SECTION("Manual HIPGraph") {
|
||||
hipGraphsManual(inputVec_h, inputVec_d, outputVec_d,
|
||||
result_d, size, maxBlocks);
|
||||
}
|
||||
|
||||
SECTION("HIPGraphs Using StreamCapture") {
|
||||
hipGraphsUsingStreamCapture(inputVec_h, inputVec_d,
|
||||
outputVec_d, result_d, size, maxBlocks);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(inputVec_d));
|
||||
HIP_CHECK(hipFree(outputVec_d));
|
||||
HIP_CHECK(hipFree(result_d));
|
||||
free(inputVec_h);
|
||||
}
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user