AI/rocm-systems
2
0
Forc 0
Cód Saincheisteanna Iarratais Tarraingthe Gníomhartha Pacáistí Tionscadail Scaoileann Vicí Gníomhaíocht

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Brabhsáil Foinse 
Change-Id: I5f8b3c3ae0eca02875c75263190d01a985077eaa
Tá an tiomantas seo le fáil i:
Jenkins
2022-03-13 19:12:18 -04:00
tuismitheoir f5c415e0d3 9a175e1654
tiomantas 242c1f892f
D'athraigh 23 comhad le 2542 breiseanna agus 262 scriosta
Íoslódáil an comhad paiste Íoslódáil Comhad Diff Leathnaigh gach comhaid Laghdaigh gach comhaid
INSTALL.md
+23 -135
Féach ar an gComhad
@@ -2,65 +2,34 @@
<!-- toc -->
- [Installing pre-built packages](#installing-pre-built-packages)
- [Install HIP](#installing-hip)
* [Prerequisites](#prerequisites)
* [AMD Platform](#amd-platform)
* [NVIDIA Platform](#nvidia-platform)
- [Building HIP from source on AMD platform](#building-hip-from-source-on-amd-platform)
* [Get HIP source code](#get-hip-source-code)
* [Set the environment variables](#set-the-environment-variables)
* [Build HIP](#build-hip)
* [Default paths and environment variables](#default-paths-and-environment-variables)
- [Building HIP from source on NVIDIA platform](#building-hip-from-source-on-NVIDIA-platform)
* [Get HIP source code](#get-hip-source-code)
* [Set the environment variables](#set-the-environment-variables)
* [Build HIP](#build-hip)
- [Verify your installation](#verify-your-installation)
- [How to build HIP from source](#how-to-build-hip-from-source)
<!-- tocstop -->
# Installing pre-built packages
# Install HIP
HIP can be easily installed using pre-built binary packages using the package manager for your platform.
HIP can be installed either on AMD ROCm platform with HIP-Clang compiler, or a CUDA platform with nvcc installed.
Note: The version definition for the HIP runtime is different from CUDA. Users can use hipRuntimeGerVersion function, on the AMD platform it returns the HIP runtime version, while on the NVIDIA platform, it returns the CUDA runtime version. There is no mapping or correlation between the HIP version and CUDA version.
## Prerequisites
HIP code can be developed either on AMD ROCm platform using HIP-Clang compiler, or a CUDA platform with nvcc installed.
On AMD platform, see Prerequisite Actions in ROCm_Installation_Guide (https://docs.amd.com/) for the release.
On NVIDIA platform, check system requirements in NVIDIA CUDA Installation Guide at https://docs.nvidia.com/cuda/cuda-installation-guide-linux/.
## 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
sudo apt install comgr
sudo apt-get -y install rocm-dkms
```
Public link for Rocm installation
https://rocmdocs.amd.com/en/latest/Installation_Guide/Installation-Guide.html
HIP-Clang is the compiler for compiling HIP programs on AMD platform.
HIP is part of ROCM packages, it will be automatically installed following the ROCm Installation Guide on AMD public documentation site (https://docs.amd.com/) for the coresponding ROCm release.
By default, HIP is installed into /opt/rocm/hip.
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=<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=<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=<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
```
## NVIDIA Platform
HIP-nvcc is the compiler for HIP program compilation on NVIDIA platform.
* Install Nvidia Driver
```
@@ -70,107 +39,26 @@ sudo reboot
Or download the latest cuda-toolkit at https://developer.nvidia.com/cuda-downloads
The driver will be installed automatically.
* Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories).
* Install the "hip-runtime-nvidia" and "hip-dev" package. This will install CUDA SDK and the HIP porting layer.
* Add the ROCm package server to your system as per the OS-specific guide in ROCm Installation Guide (https://docs.amd.com/) for the release.
* Install the "hip-runtime-nvidia" and "hip-dev" package. This will install CUDA SDK and the HIP porting layer.
```
apt-get install hip-runtime-nvidia hip-dev
```
* 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 <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.
* Default paths:
* By default HIP looks for CUDA SDK in /usr/local/cuda.
* By default HIP is installed into /opt/rocm/hip.
* Optionally, consider adding /opt/rocm/bin to your path to make it easier to use the tools.
# Building HIP from source on AMD platform
## Get HIP source code
```
git clone -b develop https://github.com/ROCm-Developer-Tools/hipamd.git
git clone -b develop https://github.com/ROCm-Developer-Tools/hip.git
git clone -b develop https://github.com/ROCm-Developer-Tools/ROCclr.git
git clone -b develop https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git
```
## Set the environment variables
```
export HIPAMD_DIR="$(readlink -f hipamd)"
export HIP_DIR="$(readlink -f hip)"
export ROCclr_DIR="$(readlink -f ROCclr)"
export OPENCL_DIR="$(readlink -f ROCm-OpenCL-Runtime)"
```
ROCclr is defined on AMD platform that HIP use Radeon Open Compute Common Language Runtime (ROCclr), which is a virtual device interface that HIP runtimes interact with different backends.
See https://github.com/ROCm-Developer-Tools/ROCclr
HIPAMD repository provides implementation specifically for AMD platform.
See https://github.com/ROCm-Developer-Tools/hipamd
## Build HIP
```
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="<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 "<ROCM_PATH>/hip".
By default, release version of AMDHIP is built.
## Default paths and environment variables
* 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
# Building HIP from source on NVIDIA platform
## Get HIP source code
```
git clone -b develop https://github.com/ROCm-Developer-Tools/hip.git
git clone -b develop https://github.com/ROCm-Developer-Tools/hipamd.git
```
## Set the environment variables
```
export HIP_DIR="$(readlink -f hip)"
export HIPAMD_DIR="$(readlink -f hipamd)"
```
## Build HIP
```
cd "$HIPAMD_DIR"
mkdir -p build; cd build
cmake -DHIP_COMMON_DIR=$HIP_DIR -DHIP_PLATFORM=nvidia -DCMAKE_INSTALL_PREFIX=$PWD/install ..
make -j$(nproc)
sudo make install
```
# Verify your installation
Run hipconfig (instructions below assume default installation path) :
```shell
<ROCM_PATH>/bin/hipconfig --full
```
or
```shell
$PWD/install/bin/hipconfig --full
/opt/rocm/bin/hipconfig --full
```
Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-4.5.x/samples/0_Intro/square).
# How to build HIP from source
Developers can build HIP from source on either AMD or NVIDIA platforms, see detail instructions at [building HIP] (docs/markdown/hip_build.md).
bin/hipcc.pl
-1
Féach ar an gComhad
@@ -697,7 +697,6 @@ if ($HIP_PLATFORM eq "amd") {
if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") {
$HIPCXXFLAGS .= " --hip-device-lib-path=\"$DEVICE_LIB_PATH\"";
}
$HIPCXXFLAGS .= " -fhip-new-launch-api";
}
if (not $isWindows) {
$HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm -lrt";
bin/hipvars.pm
+1 -1
Féach ar an gComhad
@@ -26,7 +26,7 @@ use Cwd;
use File::Basename;
$HIP_BASE_VERSION_MAJOR = "5";
$HIP_BASE_VERSION_MINOR = "1";
$HIP_BASE_VERSION_MINOR = "2";
$HIP_BASE_VERSION_PATCH = "0";
#---
docs/markdown/hip_build.md
Inrite
+137
Féach ar an gComhad
@@ -0,0 +1,137 @@
## Table of Contents
<!-- toc -->
- [Prerequisites](#Prerequisites)
- [Build HIP on AMD platform](#build-hip-on-amd-platform)
* [Get HIP source code](#get-hip-source-code)
* [Set the environment variables](#set-the-environment-variables)
* [Build HIP](#build-hip)
* [Default paths and environment variables](#default-paths-and-environment-variables)
- [Build HIP on NVIDIA platform](#build-hip-on-NVIDIA-platform)
* [Get HIP source code](#get-hip-source-code)
* [Set the environment variables](#set-the-environment-variables)
* [Build HIP](#build-hip)
- [Run HIP](#run-hip)
<!-- tocstop -->
# Prerequisites
HIP code can be developed either on AMD ROCm platform using HIP-Clang compiler, or a CUDA platform with nvcc installed.
Before build and run HIP, make sure drivers and pre-build packages are installed properly on the platform.
## AMD platform
Install ROCm packages (see ROCm Installation Guide on AMD public documentation site (https://docs.amd.com/)) or install pre-built binary packages using the package manager,
```
sudo apt install mesa-common-dev
sudo apt install clang
sudo apt install comgr
sudo apt-get -y install rocm-dkms
```
## NVIDIA platform
Install Nvidia driver and pre-build packages (see HIP Installation Guide at https://docs.amd.com/ for the release)
## Branch of repository
Before get HIP source code, set the expected branch of repository at the variable HIP_BRANCH.
For example, for ROCm5.0 release branch, set
```
export HIP_BRANCH=rocm-5.0.x
```
ROCm5.1 release branch, set
```
export HIP_BRANCH=rocm-5.1.x
```
Similiar format for future branches.
ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is at /opt/rocm.
# Build HIP on AMD platform
## Get HIP source code
```
git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hipamd.git
git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hip.git
git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/ROCclr.git
git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git
```
## Set the environment variables
```
export HIPAMD_DIR="$(readlink -f hipamd)"
export HIP_DIR="$(readlink -f hip)"
export ROCclr_DIR="$(readlink -f ROCclr)"
export OPENCL_DIR="$(readlink -f ROCm-OpenCL-Runtime)"
```
ROCclr is defined on AMD platform that HIP use Radeon Open Compute Common Language Runtime (ROCclr), which is a virtual device interface that HIP runtimes interact with different backends.
See https://github.com/ROCm-Developer-Tools/ROCclr
HIPAMD repository provides implementation specifically for AMD platform.
See https://github.com/ROCm-Developer-Tools/hipamd
## Build HIP
```
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="<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 "<ROCM_PATH>/hip".
By default, release version of AMDHIP is built.
## Default paths and environment variables
* 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 make install command, make sure HIP_PATH is pointed to $PWD/install/hip.
# Build HIP on NVIDIA platform
## Get HIP source code
```
git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hip.git
git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hipamd.git
```
## Set the environment variables
```
export HIP_DIR="$(readlink -f hip)"
export HIPAMD_DIR="$(readlink -f hipamd)"
```
## Build HIP
```
cd "$HIPAMD_DIR"
mkdir -p build; cd build
cmake -DHIP_COMMON_DIR=$HIP_DIR -DHIP_PLATFORM=nvidia -DCMAKE_INSTALL_PREFIX=$PWD/install ..
make -j$(nproc)
sudo make install
```
# Run HIP
Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-5.0.x/samples/0_Intro/square).
docs/markdown/hip_debugging.md
+51 -28
Féach ar an gComhad
@@ -131,40 +131,63 @@ There are also other debugging tools available online developers can google and
## Debugging HIP Applications
Below is an example to show how to get useful information from the debugger while running an application, which caused an issue of GPUVM fault.
Below is an example to show how to get useful information from the debugger while running a simple memory copy test, which caused an issue of segmentation fault.
```
Memory access fault by GPU node-1 on address 0x5924000. Reason: Page not present or supervisor privilege.
test: simpleTest2<?> numElements=4194304 sizeElements=4194304 bytes
Segmentation fault (core dumped)
Program received signal SIGABRT, Aborted.
[Switching to Thread 0x7fffdffb5700 (LWP 14893)]
0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
56 ../nptl/sysdeps/unix/sysv/linux/raise.c: No such file or directory.
(gdb) run
Starting program: /home/test/hipamd/build/directed_tests/runtimeApi/memory/hipMemcpy_simple
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
Breakpoint 1, main (argc=1, argv=0x7fffffffdea8)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:147
147 int main(int argc, char* argv[]) {
(gdb) c
Continuing.
[New Thread 0x7ffff64c4700 (LWP 146066)]
Thread 1 "hipMemcpy_simpl" received signal SIGSEGV, Segmentation fault.
0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:104
104 A_h1[i] = 3.14f + 1000 * i;
(gdb) bt
#0 0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
#1 0x00007ffff205b028 in __GI_abort () at abort.c:89
#2 0x00007ffff6f960eb in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#3 0x00007ffff6f99ea5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#4 0x00007ffff6f78107 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#5 0x00007ffff744f184 in start_thread (arg=0x7fffdffb5700) at pthread_create.c:312
#6 0x00007ffff211b37d in clone () at ../sysdeps/unix/sysv/linux/x86_64/clone.S:111
(gdb) info threads
Id Target Id Frame
4 Thread 0x7fffdd521700 (LWP 14895) "caffe" pthread_cond_wait@@GLIBC_2.3.2 () at ../nptl/sysdeps/unix/sysv/linux/x86_64/pthread_cond_wait.S:185
3 Thread 0x7fffddd22700 (LWP 14894) "caffe" pthread_cond_wait@@GLIBC_2.3.2 () at ../nptl/sysdeps/unix/sysv/linux/x86_64/pthread_cond_wait.S:185
* 2 Thread 0x7fffdffb5700 (LWP 14893) "caffe" 0x00007ffff2057c37 in __GI_raise (sig=sig@entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
1 Thread 0x7ffff7fa6ac0 (LWP 14892) "caffe" 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#0 0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:104
#1 0x000000000020e96c in main (argc=<optimized out>, argv=<optimized out>)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:163
(gdb) info thread
Id Target Id Frame
* 1 Thread 0x7ffff64c5880 (LWP 146060) "hipMemcpy_simpl" 0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:104
2 Thread 0x7ffff64c4700 (LWP 146066) "hipMemcpy_simpl" 0x00007ffff6b0850b in ioctl
() from /lib/x86_64-linux-gnu/libc.so.6
(gdb) thread 2
[Switching to thread 2 (Thread 0x7ffff64c4700 (LWP 146066))]
#0 0x00007ffff6b0850b in ioctl () from /lib/x86_64-linux-gnu/libc.so.6
(gdb) bt
#0 0x00007ffff6b0850b in ioctl () from /lib/x86_64-linux-gnu/libc.so.6
#1 0x00007ffff6604568 in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#2 0x00007ffff65fe73a in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#3 0x00007ffff659e4d6 in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#4 0x00007ffff65807de in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#5 0x00007ffff65932a2 in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#6 0x00007ffff654f547 in ?? () from /opt/rocm/lib/libhsa-runtime64.so.1
#7 0x00007ffff7f76609 in start_thread () from /lib/x86_64-linux-gnu/libpthread.so.0
#8 0x00007ffff6b13293 in clone () from /lib/x86_64-linux-gnu/libc.so.6
(gdb) thread 1
[Switching to thread 1 (Thread 0x7ffff7fa6ac0 (LWP 14892))]
#0 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
[Switching to thread 1 (Thread 0x7ffff64c5880 (LWP 146060))]
#0 0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:104
104 A_h1[i] = 3.14f + 1000 * i;
(gdb) bt
#0 0x00007ffff6f934d5 in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#1 0x00007ffff6f929ba in ?? () from /opt/rocm/hsa/lib/libhsa-runtime64.so.1
#2 0x00007fffe080beca in HSADispatch::waitComplete() () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#3 0x00007fffe080415f in HSADispatch::dispatchKernelAsync(Kalmar::HSAQueue*, void const*, int, bool) () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#4 0x00007fffe080238e in Kalmar::HSAQueue::dispatch_hsa_kernel(hsa_kernel_dispatch_packet_s const*, void const*, unsigned long, hc::completion_future*) () from /opt/rocm/hcc/lib/libmcwamp_hsa.so
#5 0x00007ffff7bb7559 in hipModuleLaunchKernel () from /opt/rocm/hip/lib/libhip_hcc.so
#6 0x00007ffff2e6cd2c in mlopen::HIPOCKernel::run (this=0x7fffffffb5a8, args=0x7fffffffb2a8, size=80) at /root/MIOpen/src/hipoc/hipoc_kernel.cpp:15
#0 0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:104
#1 0x000000000020e96c in main (argc=<optimized out>, argv=<optimized out>)
at /home/test/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp:163
(gdb)
...
```
docs/markdown/hip_kernel_language.md
+7 -4
Féach ar an gComhad
@@ -125,6 +125,8 @@ MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
The hipLaunchKernelGGL macro always starts with the five parameters specified above, followed by the kernel arguments. HIPIFY tools optionally convert Cuda launch syntax to hipLaunchKernelGGL, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernelGGL parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (thread*, block*, grid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins).
Please note, HIP does not support kernel launch with total work items defined in dimension with size gridDim x blockDim >= 2^32.
## Kernel-Launch Example
```
@@ -183,6 +185,7 @@ The `__restrict__` keyword tells the compiler that the associated memory pointer
Built-ins determine the coordinate of the active work item in the execution grid. They are defined in amd_hip_runtime.h (rather than being implicitly defined by the compiler).
In HIP, built-ins coordinate variable definitions are the same as in Cuda, for instance:
threadIdx.x, blockIdx.y, gridDim.y, etc.
The products gridDim.x * blockDim.x, gridDim.y * blockDim.y and gridDim.z * blockDim.z are always less than 2^32.
### warpSize
The warpSize variable is of type int and contains the warp size (in threads) for the target device. Note that all current Nvidia devices return 32 for this variable, and all current AMD devices return 64. Device code should use the warpSize built-in to develop portable wave-aware code.
@@ -214,12 +217,12 @@ HIP supports the following short vector formats:
- double1, double2, double3, double4
### dim3
dim3 is a three-dimensional integer vector type commonly used to specify grid and group dimensions. Unspecified dimensions are initialized to 1.
dim3 is a three-dimensional integer vector type commonly used to specify grid and group dimensions. Unspecified dimensions are initialized to 1.
```
typedef struct dim3 {
uint32_t x;
uint32_t y;
uint32_t z;
uint32_t x;
uint32_t y;
uint32_t z;
dim3(uint32_t _x=1, uint32_t _y=1, uint32_t _z=1) : x(_x), y(_y), z(_z) {};
};
include/hip/hip_bfloat16.h
+1 -1
Féach ar an gComhad
@@ -56,7 +56,7 @@ struct hip_bfloat16
enum truncate_t
{
truncate_0
truncate
};
__host__ __device__ hip_bfloat16() = default;
include/hip/hip_runtime_api.h
+257 -68
Féach ar an gComhad
@@ -2680,8 +2680,8 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
* If no variable of that name exists, it returns hipErrorNotFound. Both parameters dptr and bytes are optional.
* If one of them is NULL, it is ignored and hipSuccess is returned.
*
* @param[out] dptr Returned global device pointer
* @param[out] bytes Returned global size in bytes
* @param[out] dptr Returns global device pointer
* @param[out] bytes Returns global size in bytes
* @param[in] hmod Module to retrieve global from
* @param[in] name Name of global to retrieve
*
@@ -2690,17 +2690,97 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
*/
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
hipModule_t hmod, const char* name);
/**
* @brief Gets device pointer associated with symbol on the device.
*
* @param[out] devPtr pointer to the device associated the symbole
* @param[in] symbol pointer to the symbole of the device
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol);
/**
* @brief Gets the size of the given symbol on the device.
*
* @param[in] symbol pointer to the device symbole
* @param[out] size pointer to the size
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGetSymbolSize(size_t* size, const void* symbol);
/**
* @brief Copies data to the given symbol on the device.
* Symbol HIP APIs allow a kernel to define a device-side data symbol which can be accessed on
* the host side. The symbol can be in __constant or device space.
* Note that the symbol name needs to be encased in the HIP_SYMBOL macro.
* This also applies to hipMemcpyFromSymbol, hipGetSymbolAddress, and hipGetSymbolSize.
* For detail usage, see the example at
* https://github.com/ROCm-Developer-Tools/HIP/blob/rocm-5.0.x/docs/markdown/hip_porting_guide.md
*
* @param[out] symbol pointer to the device symbole
* @param[in] src pointer to the source address
* @param[in] sizeBytes size in bytes to copy
* @param[in] offset offset in bytes from start of symbole
* @param[in] kind type of memory transfer
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipMemcpyToSymbol(const void* symbol, const void* src,
size_t sizeBytes, size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
/**
* @brief Copies data to the given symbol on the device asynchronously.
*
* @param[out] symbol pointer to the device symbole
* @param[in] src pointer to the source address
* @param[in] sizeBytes size in bytes to copy
* @param[in] offset offset in bytes from start of symbole
* @param[in] kind type of memory transfer
* @param[in] stream stream identifier
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src,
size_t sizeBytes, size_t offset,
hipMemcpyKind kind, hipStream_t stream __dparm(0));
/**
* @brief Copies data from the given symbol on the device.
*
* @param[out] dptr Returns pointer to destinition memory address
* @param[in] symbol pointer to the symbole address on the device
* @param[in] sizeBytes size in bytes to copy
* @param[in] offset offset in bytes from the start of symbole
* @param[in] kind type of memory transfer
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol,
size_t sizeBytes, size_t offset __dparm(0),
hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
/**
* @brief Copies data from the given symbol on the device asynchronously.
*
* @param[out] dptr Returns pointer to destinition memory address
* @param[in] symbol pointer to the symbole address on the device
* @param[in] sizeBytes size in bytes to copy
* @param[in] offset offset in bytes from the start of symbole
* @param[in] kind type of memory transfer
* @param[in] stream stream identifier
*
* @return #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol,
size_t sizeBytes, size_t offset,
hipMemcpyKind kind,
@@ -3747,6 +3827,10 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
* must be in the memory layout and alignment expected by the kernel.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32. So gridDim.x * blockDim.x, gridDim.y * blockDim.y
* and gridDim.z * blockDim.z are always less than 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
@@ -3770,6 +3854,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th
* default stream is used with associated synchronization rules.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue, hipErrorCooperativeLaunchTooLarge
*/
hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX,
@@ -3820,6 +3907,9 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
* @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue
*/
//TODO - Match CUoccupancyB2DSize
@@ -3836,6 +3926,9 @@ hipError_t hipModuleOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize
* @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit
* @param [in] flags Extra flags for occupancy calculation (only default supported)
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue
*/
//TODO - Match CUoccupancyB2DSize
@@ -3893,6 +3986,9 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
* @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue
*/
hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
@@ -3948,6 +4044,9 @@ hipError_t hipProfilerStop();
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case the
* default stream is used with associated synchronization rules.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
*/
@@ -3982,6 +4081,9 @@ hipError_t hipLaunchByPtr(const void* func);
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case the
* default stream is used with associated synchronization rules.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
*/
@@ -3999,6 +4101,12 @@ hipError_t __hipPushCallConfiguration(dim3 gridDim,
* @param [out] stream Stream where the kernel should be dispatched. May be 0, in which case the
* default stream is used with associated synchronization rules.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
*/
@@ -4070,6 +4178,152 @@ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3
* @{
* This section describes the texture management functions of HIP runtime API.
*/
/**
* @brief Binds a mipmapped array to a texture.
*
* @param [in] tex pointer to the texture reference to bind
* @param [in] mipmappedArray memory mipmapped array on the device
* @param [in] desc opointer to the channel format
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipBindTextureToMipmappedArray(
const textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc);
/**
* @brief Gets the texture reference related with the symbol.
*
* @param [out] texref texture reference
* @param [in] symbol pointer to the symbol related with the texture for the reference
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipGetTextureReference(
const textureReference** texref,
const void* symbol);
/**
* @brief Creates a texture object.
*
* @param [out] pTexObject pointer to the texture object to create
* @param [in] pResDesc pointer to resource descriptor
* @param [in] pTexDesc pointer to texture descriptor
* @param [in] pResViewDesc pointer to resource view descriptor
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipCreateTextureObject(
hipTextureObject_t* pTexObject,
const hipResourceDesc* pResDesc,
const hipTextureDesc* pTexDesc,
const struct hipResourceViewDesc* pResViewDesc);
/**
* @brief Destroys a texture object.
*
* @param [in] textureObject texture object to destroy
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
/**
* @brief Gets the channel descriptor in an array.
*
* @param [in] desc pointer to channel format descriptor
* @param [out] array memory array on the device
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipGetChannelDesc(
hipChannelFormatDesc* desc,
hipArray_const_t array);
/**
* @brief Gets resource descriptor for the texture object.
*
* @param [out] pResDesc pointer to resource descriptor
* @param [in] textureObject texture object
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipGetTextureObjectResourceDesc(
hipResourceDesc* pResDesc,
hipTextureObject_t textureObject);
/**
* @brief Gets resource view descriptor for the texture object.
*
* @param [out] pResViewDesc pointer to resource view descriptor
* @param [in] textureObject texture object
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
hipError_t hipGetTextureObjectResourceViewDesc(
struct hipResourceViewDesc* pResViewDesc,
hipTextureObject_t textureObject);
/**
* @brief Gets texture descriptor for the texture object.
*
* @param [out] pTexDesc pointer to texture descriptor
* @param [in] textureObject texture object
*
* @returns hipSuccess, hipErrorInvalidValue
*
*/
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 TextureD Texture Management [Deprecated]
@@ -4167,73 +4421,7 @@ hipError_t hipTexRefSetMaxAnisotropy(
/**
* @}
*/
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);
// doxygen end Texture management
/**
* @addtogroup TextureU Texture Management [Unsupported]
* @{
* @ingroup Texture
* This section describes the unsupported texture management functions of HIP runtime API.
*/
// The following are not supported.
/**
@@ -4275,6 +4463,7 @@ hipError_t hipMipmappedArrayGetLevel(
/**
* @}
*/
// doxygen end Texture management
/**
* @}
tests/catch/CMakeLists.txt
+14 -11
Féach ar an gComhad
@@ -85,7 +85,7 @@ set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake)
if(HIP_PLATFORM MATCHES "amd" AND HIP_COMPILER MATCHES "clang")
add_compile_options(-Wall -Wextra -pedantic -Werror)
add_compile_options(-Wall -Wextra -pedantic -Werror -Wno-deprecated)
endif()
cmake_policy(PUSH)
@@ -109,16 +109,19 @@ if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerat
RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT)
# Trim out gfx000
string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH})
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
if (NOT HIP_GPU_ARCH STREQUAL "")
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
endif()
else()
message(STATUS "ROCm Agent Enumurator found no valid architectures")
endif()
endif()
tests/catch/unit/CMakeLists.txt
+2 -2
Féach ar an gComhad
@@ -18,10 +18,11 @@
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
# Currently build fails for these two test sets
# build fails in windows
if (UNIX)
add_subdirectory(memory)
add_subdirectory(graph)
add_subdirectory(rtc)
endif()
add_subdirectory(deviceLib)
@@ -29,7 +30,6 @@ add_subdirectory(stream)
add_subdirectory(event)
add_subdirectory(occupancy)
add_subdirectory(device)
add_subdirectory(rtc)
add_subdirectory(printf)
add_subdirectory(printfExe)
add_subdirectory(texture)
tests/catch/unit/cooperativeGrps/coalesced_groups_shfl_down.cc
+265
Féach ar an gComhad
@@ -0,0 +1,265 @@
/*
Copyright (c) 2020 - 2022 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 Description:
/* This test implements sum reduction kernel, first with each threads own rank
as input and comparing the sum with expected sum output derieved from n(n-1)/2
formula.
This sample tests functionality of intrinsics provided by thread_block_tile type,
shfl_down and shfl_xor.
*/
#include <hip_test_common.hh>
#include <hip/hip_cooperative_groups.h>
#include <stdio.h>
#include <vector>
using namespace cooperative_groups;
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
#define WAVE_SIZE 32
__device__ int reduction_kernel_shfl_down(coalesced_group const& g, int val) {
int sz = g.size();
for (int i = sz / 2; i > 0; i >>= 1) {
val += g.shfl_down(val, i);
}
// Choose the 0'th indexed thread that holds the reduction value to return
if (g.thread_rank() == 0) {
return val;
}
// Rest of the threads return no useful values
else {
return -1;
}
}
__global__ void kernel_shfl_down (int * dPtr, int *dResults, int lane_delta, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group const& g = coalesced_threads();
int rank = g.thread_rank();
int val = dPtr[rank];
dResults[rank] = g.shfl_down(val, lane_delta);
return;
}
}
__global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group threadBlockCGTy = coalesced_threads();
int input, outputSum, expectedSum;
// Choose a leader thread to print the results
if (threadBlockCGTy.thread_rank() == 0) {
printf(" Creating %d groups, of tile size %d threads:\n\n",
(int)threadBlockCGTy.size() / tileSz, tileSz);
}
threadBlockCGTy.sync();
coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
int threadRank = tiledPartition.thread_rank();
input = tiledPartition.thread_rank();
// (n-1)(n)/2
expectedSum = ((tileSz - 1) * tileSz / 2);
outputSum = reduction_kernel_shfl_down(tiledPartition, input);
if (tiledPartition.thread_rank() == 0) {
printf(
" Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected "
"%d)\n",
tiledPartition.size() - 1, outputSum, expectedSum);
result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum;
}
return;
}
}
void verifyResults(int* ptr, int expectedResult, int numTiles) {
for (int i = 0; i < numTiles; i++) {
if (ptr[i] != expectedResult) {
INFO(" Results do not match! ");
}
}
}
void compareResults(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" results do not match.");
}
}
}
void printResults(int* ptr, int size) {
for (int i = 0; i < size; i++) {
std::cout << ptr[i] << " ";
}
std::cout << '\n';
}
static void test_group_partition(unsigned int tileSz) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = 32;
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz;
int expectedSum = ((tileSz - 1) * tileSz / 2);
int* expectedResult = new int[numTiles];
// numTiles = 0 when partitioning is possible. The below statement is to avoid
// out-of-bounds error and still evaluate failure case.
numTiles = (numTiles == 0) ? 1 : numTiles;
for (int i = 0; i < numTiles; i++) {
expectedResult[i] = expectedSum;
}
int* dResult = NULL;
int* hResult = NULL;
hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
memset(hResult, 0, numTiles * sizeof(int));
hipMalloc(&dResult, numTiles * sizeof(int));
// Launch Kernel
hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dResult, tileSz, i);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost);
verifyResults(hResult, expectedSum, numTiles);
// Free all allocated memory on host and device
hipFree(dResult);
hipFree(hResult);
delete[] expectedResult;
printf("\n...PASSED.\n\n");
}
}
static void test_shfl_down() {
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
int totalThreads = blockSize * threadsPerBlock;
int group_size = totalThreads / i;
int group_size_in_bytes = group_size * sizeof(int);
int* hPtr = NULL;
int* dPtr = NULL;
int* dResults = NULL;
int lane_delta = rand() % group_size;
std::cout << "Testing coalesced_groups shfl_down with lane_delta " << lane_delta << "and group size "
<< WAVE_SIZE << '\n' << std::endl;
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
}
int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
for (int i = 0; i < group_size; i++) {
cpuResultsArr[i] = (i + lane_delta >= group_size) ? hPtr[i] : hPtr[i + lane_delta];
}
//printf("Array passed to GPU for computation\n");
//printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
//printf("GPU results: \n");
//printResults(hPtr, WAVE_SIZE);
//printf("Printing cpu to be verified array\n");
//printResults(cpuResultsArr, WAVE_SIZE);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
free(cpuResultsArr);
}
}
TEST_CASE("Unit_coalesced_groups_shfl_down") {
// Use default device for validating the test
int deviceId;
ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
hipDeviceProp_t deviceProperties;
ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
// Test shfl_down with random group sizes
for (int i = 0; i < 100; i++) {
test_shfl_down();
}
std::cout << "Testing static tiled_partition for different tile sizes using shfl_down"
<< std::endl;
int testNo = 1;
std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
for (auto i : tileSizes) {
std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
test_group_partition(i);
testNo++;
}
}
tests/catch/unit/cooperativeGrps/coalesced_groups_shfl_up.cc
+249
Féach ar an gComhad
@@ -0,0 +1,249 @@
/*
Copyright (c) 2022 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 Description:
/* This test implements prefix sum(scan) kernel, first with each threads own rank
as input and comparing the sum with expected serial summation output on CPU.
This sample tests functionality of intrinsics provided by coalesced_group,
shfl_up.
*/
#include <hip_test_common.hh>
#include <hip/hip_cooperative_groups.h>
#include <stdio.h>
#include <vector>
using namespace cooperative_groups;
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
#define WAVE_SIZE 32
__device__ int prefix_sum_kernel(coalesced_group const& g, int val) {
int sz = g.size();
for (int i = 1; i < sz; i <<= 1) {
int temp = g.shfl_up(val, i);
if (g.thread_rank() >= i) {
val += temp;
}
}
return val;
}
__global__ void kernel_shfl_up (int * dPtr, int *dResults, int lane_delta, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group g = coalesced_threads();
int rank = g.thread_rank();
int val = dPtr[rank];
dResults[rank] = g.shfl_up(val, lane_delta);
return;
}
}
__global__ void kernel_cg_group_partition(int* dPtr, unsigned int tileSz, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group threadBlockCGTy = coalesced_threads();
int input, outputSum;
// we pass its own thread rank as inputs
input = threadBlockCGTy.thread_rank();
// Choose a leader thread to print the results
if (threadBlockCGTy.thread_rank() == 0) {
printf(" Creating %d groups, of tile size %d threads:\n\n",
(int)threadBlockCGTy.size() / tileSz, tileSz);
}
threadBlockCGTy.sync();
coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
input = tiledPartition.thread_rank();
outputSum = prefix_sum_kernel(tiledPartition, input);
// Update the result array with the corresponsing prefix sum
dPtr[threadBlockCGTy.thread_rank()] = outputSum;
return;
}
}
void serialScan(int* ptr, int size) {
// Fill up the array
for (int i = 0; i < size; i++) {
ptr[i] = i;
}
int acc = 0;
for (int i = 0; i < size; i++) {
acc = acc + ptr[i];
ptr[i] = acc;
}
}
void printResults(int* ptr, int size) {
for (int i = 0; i < size; i++) {
std::cout << ptr[i] << " ";
}
std::cout << '\n';
}
void verifyResults(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" Results do not match.");
}
}
}
static void test_group_partition(unsigned tileSz) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
int* hPtr = NULL;
int* dPtr = NULL;
int* cpuPrefixSum = NULL;
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
hipMalloc(&dPtr, arrSize);
// Launch Kernel
hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i);
hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
cpuPrefixSum = new int[tileSz];
serialScan(cpuPrefixSum, tileSz);
//std::cout << "\nPrefix sum results on CPU\n";
//printResults(cpuPrefixSum, tileSz);
//std::cout << "\nPrefix sum results on GPU\n";
//printResults(hPtr, tileSz);
std::cout << "\n";
verifyResults(hPtr, cpuPrefixSum, tileSz);
std::cout << "Results verified!\n";
delete[] cpuPrefixSum;
hipFree(hPtr);
hipFree(dPtr);
}
}
static void test_shfl_up() {
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
int totalThreads = blockSize * threadsPerBlock;
int group_size = totalThreads / i;
int group_size_in_bytes = group_size * sizeof(int);
int* hPtr = NULL;
int* dPtr = NULL;
int* dResults = NULL;
int lane_delta = (rand() % group_size);
std::cout << "Testing coalesced_groups shfl_up with lane_delta " << lane_delta
<< " and group size " << WAVE_SIZE << '\n' << std::endl;
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
}
//printResults(hPtr, WAVE_SIZE);
int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
for (int i = 0; i < group_size; i++) {
cpuResultsArr[i] = (i <= (lane_delta - 1)) ? hPtr[i] : hPtr[i - lane_delta];
}
//printf("Printing cpu results arr\n");
//printResults(cpuResultsArr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i);
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
//printf("GPU computation array :\n");
//printResults(hPtr, WAVE_SIZE);
verifyResults(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
free(cpuResultsArr);
}
}
TEST_CASE("Unit_coalesced_groups_shfl_down") {
// Use default device for validating the test
int deviceId;
ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
hipDeviceProp_t deviceProperties;
ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
for (int i = 0; i < 100; i++) {
test_shfl_up();
}
std::cout << "Testing coalesced_groups partitioning and shfl_up" << '\n' << std::endl;
int testNo = 1;
std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
for (auto i : tileSizes) {
std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
test_group_partition(i);
testNo++;
}
}
/* Kogge-Stone algorithm */
tests/catch/unit/cooperativeGrps/simple_coalesced_groups.cc
+571
Féach ar an gComhad
@@ -0,0 +1,571 @@
/*
Copyright (c) 2020 - 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 Description:
/* This test implements sum reduction kernel, first with each threads own rank
as input and comparing the sum with expected sum output derieved from n(n-1)/2
formula. The second part, partitions this parent group into child subgroups
a.k.a tiles using using tiled_partition() collective operation. This can be called
with a static tile size, passed in templated non-type variable-tiled_partition<tileSz>,
or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
cases.
This test tests functionality of cg group partitioning, (static and dynamic) and its respective
API's size(), thread_rank(), and sync().
*/
#include <hip_test_common.hh>
#include <hip/hip_cooperative_groups.h>
#include <stdio.h>
#include <vector>
using namespace cooperative_groups;
#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
#define NUM_ELEMS 10000000
#define NUM_THREADS_PER_BLOCK 512
#define WAVE_SIZE 32
/* Test coalesced group's functionality.
*
*/
__device__ int atomicAggInc(int *ptr) {
coalesced_group g = coalesced_threads();
int prev;
// elect the first active thread to perform atomic add
if (g.thread_rank() == 0) {
prev = atomicAdd(ptr, g.size());
}
// broadcast previous value within the warp
// and add each active thread’s rank to it
prev = g.thread_rank() + g.shfl(prev, 0);
return prev;
}
__global__ void kernel_shfl (int * dPtr, int *dResults, int srcLane, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group const& g = coalesced_threads();
int rank = g.thread_rank();
int val = dPtr[rank];
dResults[rank] = g.shfl(val, srcLane);
return;
}
}
__global__ void kernel_shfl_any_to_any (int *randVal, int *dsrcArr, int *dResults, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group const& g = coalesced_threads();
int rank = g.thread_rank();
int val = randVal[rank];
dResults[rank] = g.shfl(val, dsrcArr[rank]);
return;
}
}
__global__ void filter_arr(int *dst, int *nres, const int *src, int n) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = id; i < n; i += gridDim.x * blockDim.x) {
if (src[i] > 0) dst[atomicAggInc(nres)] = src[i];
}
}
/* Parallel reduce kernel.
*
* Step complexity: O(log n)
* Work complexity: O(n)
*
* Note: This kernel works only with power of 2 input arrays.
*/
__device__ int reduction_kernel(coalesced_group g, int* x, int val) {
int lane = g.thread_rank();
int sz = g.size();
for (int i = g.size() / 2; i > 0; i /= 2) {
// use lds to store the temporary result
x[lane] = val;
// Ensure all the stores are completed.
g.sync();
if (lane < i) {
val += x[lane + i];
}
// It must work on one tiled thread group at a time,
// and it must make sure all memory operations are
// completed before moving to the next stride.
// sync() here just does that.
g.sync();
}
// Choose the 0'th indexed thread that holds the reduction value to return
if (g.thread_rank() == 0) {
return val;
}
// Rest of the threads return no useful values
else {
return -1;
}
}
__global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* result,
bool isGlobalMem, int* globalMem, int cg_sizes) {
int id = threadIdx.x + blockIdx.x * blockDim.x;
if (id % cg_sizes == 0) {
coalesced_group threadBlockCGTy = coalesced_threads();
int threadBlockGroupSize = threadBlockCGTy.size();
int* workspace = NULL;
if (isGlobalMem) {
workspace = globalMem;
} else {
// Declare a shared memory
extern __shared__ int sharedMem[];
workspace = sharedMem;
}
int input, outputSum, expectedOutput;
// input to reduction, for each thread, is its' rank in the group
input = threadBlockCGTy.thread_rank();
expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
if (threadBlockCGTy.thread_rank() == 0) {
printf(" Sum of all ranks 0..%d in coalesced_group is %d\n\n",
(int)threadBlockCGTy.size() - 1, outputSum);
printf(" Creating %d groups, of tile size %d threads:\n\n",
(int)threadBlockCGTy.size() / tileSz, tileSz);
}
threadBlockCGTy.sync();
coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
// This offset allows each group to have its own unique area in the workspace array
int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
if (tiledPartition.thread_rank() == 0) {
printf(
" Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
"rank: %d\n",
tiledPartition.size() - 1, outputSum, input);
result[input / (tileSz)] = outputSum;
}
return;
}
}
__global__ void kernel_coalesced_active_groups() {
thread_block threadBlockCGTy = this_thread_block();
int threadBlockGroupSize = threadBlockCGTy.size();
// input to reduction, for each thread, is its' rank in the group
int input = threadBlockCGTy.thread_rank();
if (threadBlockCGTy.thread_rank() == 0) {
printf(" Creating odd and even set of active thread groups based on branch divergence\n\n");
}
threadBlockCGTy.sync();
// Group all active odd threads
if (threadBlockCGTy.thread_rank() % 2) {
coalesced_group activeOdd = coalesced_threads();
if (activeOdd.thread_rank() == 0) {
printf(" ODD: Size of odd set of active threads is %d."
" Corresponding parent thread_rank is %d.\n\n",
activeOdd.size(), threadBlockCGTy.thread_rank());
}
}
else { // Group all active even threads
coalesced_group activeEven = coalesced_threads();
if (activeEven.thread_rank() == 0) {
printf(" EVEN: Size of even set of active threads is %d."
" Corresponding parent thread_rank is %d.",
activeEven.size(), threadBlockCGTy.thread_rank());
}
}
return;
}
void printResults(int* ptr, int size) {
for (int i = 0; i < size; i++) {
std::cout << ptr[i] << " ";
}
std::cout << '\n';
}
void compareResults(int* cpu, int* gpu, int size) {
for (unsigned int i = 0; i < size / sizeof(int); i++) {
if (cpu[i] != gpu[i]) {
INFO(" results do not match.");
}
}
}
static void test_active_threads_grouping() {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
// Launch Kernel
hipLaunchKernelGGL(kernel_coalesced_active_groups, blockSize, threadsPerBlock, 0, 0);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
printf("\n...PASSED.\n\n");
}
// Search if the sum exists in the expected results array
void verifyResults(int* hPtr, int* dPtr, int size) {
int i = 0, j = 0;
for (i = 0; i < size; i++) {
for (j = 0; j < size; j++) {
if (hPtr[i] == dPtr[j]) {
break;
}
}
if (j == size) {
INFO(" Result verification failed!");
}
}
}
static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz;
// numTiles = 0 when partitioning is possible. The below statement is to avoid
// out-of-bounds error and still evaluate failure case.
numTiles = (numTiles == 0) ? 1 : numTiles;
// Build an array of expected reduction sum output on the host
// based on the sum of their respective thread ranks to use for verification
int* expectedSum = new int[numTiles];
int temp = 0, sum = 0;
for (int i = 1; i <= numTiles; i++) {
sum = temp;
temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
expectedSum[i-1] = temp - sum;
}
int* dResult = NULL;
hipMalloc(&dResult, sizeof(int) * numTiles);
int* globalMem = NULL;
if (useGlobalMem) {
hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
}
int* hResult = NULL;
hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
memset(hResult, 0, numTiles * sizeof(int));
// Launch Kernel
if (useGlobalMem) {
hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, 0, 0, tileSz,
dResult, useGlobalMem, globalMem, i);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
} else {
hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem, i);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
}
hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
verifyResults(expectedSum, hResult, numTiles);
// Free all allocated memory on host and device
hipFree(dResult);
hipFree(hResult);
if (useGlobalMem) {
hipFree(globalMem);
}
delete[] expectedSum;
printf("\n...PASSED.\n\n");
}
}
static void test_shfl_any_to_any() {
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
int totalThreads = blockSize * threadsPerBlock;
int group_size = (totalThreads + i - 1) / i;
int group_size_in_bytes = group_size * sizeof(int);
int* hPtr = NULL;
int* dPtr = NULL;
int* dsrcArr = NULL;
int* dResults = NULL;
int* srcArr = (int*)malloc(group_size_in_bytes);
int* srcArrCpu = (int*)malloc(group_size_in_bytes);
std::cout << "Testing coalesced_groups shfl any-to-any\n" <<std::endl;
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
}
// Fill up the random array
for (int i = 0; i < group_size; i++) {
srcArr[i] = rand() % 1000;
srcArrCpu[i] = srcArr[i] % group_size;
}
/* Fill cpu results array so that we can verify with gpu computation */
int* cpuResultsArr = (int*)malloc(group_size_in_bytes);
for(int i = 0; i < group_size; i++) {
cpuResultsArr[i] = hPtr[srcArrCpu[i]];
}
//printf("Array passed to GPU for computation\n");
//printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
hipMalloc(&dsrcArr, group_size_in_bytes);
hipMemcpy(dsrcArr, srcArr, group_size_in_bytes, hipMemcpyHostToDevice);
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl_any_to_any, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0 , dPtr, dsrcArr, dResults, i);
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
//printf("GPU results: \n");
//printResults(hPtr, group_size);
//printf("Printing cpu to be verified array\n");
//printResults(cpuResultsArr, group_size);
//printf("Printing srcLane array that was passed\n");
//printResults(srcArr, group_size);
//printf("Printing srcLane array on the CPU\n");
//printResults(srcArrCpu, group_size);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
free(srcArr);
free(srcArrCpu);
free(cpuResultsArr);
}
}
static void test_shfl_broadcast() {
std::vector<unsigned int> cg_sizes = {1, 2, 3};
for (auto i : cg_sizes) {
hipError_t err;
int blockSize = 1;
int threadsPerBlock = WAVE_SIZE;
int totalThreads = blockSize * threadsPerBlock;
int group_size = (totalThreads + i - 1) / i;
int group_size_in_bytes = group_size * sizeof(int);
int* hPtr = NULL;
int* dPtr = NULL;
int* dResults = NULL;
int srcLane = rand() % 1000;
int srcLaneCpu = 0;
std::cout << "Testing coalesced_groups shfl with srcLane " << srcLane << '\n'
<< " and group size " << i <<std::endl;
int arrSize = blockSize * threadsPerBlock * sizeof(int);
hipHostMalloc(&hPtr, arrSize);
// Fill up the array
for (int i = 0; i < WAVE_SIZE; i++) {
hPtr[i] = rand() % 1000;
}
/* Fill cpu results array so that we can verify with gpu computation */
srcLaneCpu = hPtr[srcLane % group_size];
int* cpuResultsArr = (int*)malloc(sizeof(int) * group_size);
for (int i = 0; i < group_size; i++) {
cpuResultsArr[i] = srcLaneCpu;
}
printf("Array passed to GPU for computation\n");
printResults(hPtr, WAVE_SIZE);
hipMalloc(&dPtr, group_size_in_bytes);
hipMalloc(&dResults, group_size_in_bytes);
hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice);
// Launch Kernel
hipLaunchKernelGGL(kernel_shfl, blockSize, threadsPerBlock,
threadsPerBlock * sizeof(int), 0, dPtr, dResults, srcLane, i);
hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost);
err = hipDeviceSynchronize();
if (err != hipSuccess) {
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
}
printf("GPU results: \n");
printResults(hPtr, group_size);
printf("Printing cpu to be verified array\n");
printResults(cpuResultsArr, group_size);
compareResults(hPtr, cpuResultsArr, group_size_in_bytes);
std::cout << "Results verified!\n";
hipFree(hPtr);
hipFree(dPtr);
free(cpuResultsArr);
}
}
TEST_CASE("Unit_coalesced_groups") {
// Use default device for validating the test
int deviceId;
HIP_CHECK(hipGetDevice(&deviceId));
hipDeviceProp_t deviceProperties;
HIP_CHECK(hipGetDeviceProperties(&deviceProperties, deviceId));
int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
std::cout << "Now testing coalesced_groups" << '\n' << std::endl;
int *data_to_filter, *filtered_data, nres = 0;
int *d_data_to_filter, *d_filtered_data, *d_nres;
int numOfBuckets = 5;
data_to_filter = reinterpret_cast<int *>(malloc(sizeof(int) * NUM_ELEMS));
// Generate input data.
for (int i = 0; i < NUM_ELEMS; i++) {
data_to_filter[i] = rand() % numOfBuckets;
}
HIP_CHECK(hipMalloc(&d_data_to_filter, sizeof(int) * NUM_ELEMS));
HIP_CHECK(hipMalloc(&d_filtered_data, sizeof(int) * NUM_ELEMS));
HIP_CHECK(hipMalloc(&d_nres, sizeof(int)));
HIP_CHECK(hipMemcpy(d_data_to_filter, data_to_filter,
sizeof(int) * NUM_ELEMS, hipMemcpyHostToDevice));
hipMemset(d_nres, 0, sizeof(int));
dim3 dimBlock(NUM_THREADS_PER_BLOCK, 1, 1);
dim3 dimGrid((NUM_ELEMS / NUM_THREADS_PER_BLOCK) + 1, 1, 1);
filter_arr<<<dimGrid, dimBlock>>>(d_filtered_data, d_nres, d_data_to_filter,
NUM_ELEMS);
HIP_CHECK(hipMemcpy(&nres, d_nres, sizeof(int), hipMemcpyDeviceToHost));
filtered_data = reinterpret_cast<int *>(malloc(sizeof(int) * nres));
HIP_CHECK(hipMemcpy(filtered_data, d_filtered_data, sizeof(int) * nres,
hipMemcpyDeviceToHost));
int *host_filtered_data =
reinterpret_cast<int *>(malloc(sizeof(int) * NUM_ELEMS));
// Generate host output with host filtering code.
int host_flt_count = 0;
for (int i = 0; i < NUM_ELEMS; i++) {
if (data_to_filter[i] > 0) {
host_filtered_data[host_flt_count++] = data_to_filter[i];
}
}
printf("\nWarp Aggregated Atomics %s \n",
(host_flt_count == nres) ? "PASSED" : "FAILED");
// Now, testing shfl collective
std::cout << "Now testing shfl collective as a broadcast" << '\n' << std::endl;
for (int i = 0; i < 100; i++) {
test_shfl_broadcast();
}
// Now, testing shfl collective
std::cout << "Now testing shfl operations any-to-any member lanes" << '\n' << std::endl;
for (int i = 0; i < 100; i++) {
test_shfl_any_to_any();
}
// Now, pass a already coalesced_group that was partitioned
/* Test coalesced group partitioning */
std::cout << "Now testing coalesced_groups partitioning" << '\n' << std::endl;
int testNo = 1;
for (int memTy = 0; memTy < 2; memTy++) {
std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
for (auto i : tileSizes) {
std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
test_group_partition(i, memTy);
testNo++;
}
}
std::cout << "Now grouping active threads based on branch divergence" << '\n' << std::endl;
test_active_threads_grouping();
}
tests/catch/unit/graph/CMakeLists.txt
+1
Féach ar an gComhad
@@ -41,6 +41,7 @@ set(TEST_SRC
hipGraphHostNodeSetParams.cc
hipGraphAddMemcpyNode1D.cc
hipGraphAddChildGraphNode.cc
hipGraphNodeGetType.cc
)
hip_add_exe_to_target(NAME GraphsTest
tests/catch/unit/graph/hipGraphNodeGetType.cc
+195
Féach ar an gComhad
@@ -0,0 +1,195 @@
/*
Copyright (c) 2022 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 :
Functional ::
1) Pass different types of valid nodes and get the corresponding type of the node
2) Add Graph node, destroy the graph node and add the same node of different type.
hipGraphNodeGetTye should return the new type
3) Add graph node, overwrite new type to the same graph node and trigger
hipGraphNodeGetType which should return the updated type.
Negative ::
1) Pass nullptr to graph node
2) Pass nullptr to node type
3) Pass invalid to graph node
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
TEST_CASE("Unit_hipGraphNodeGetType_Negative") {
SECTION("Pass nullptr to graph node") {
hipGraphNodeType nodeType;
REQUIRE(hipGraphNodeGetType(nullptr, &nodeType) == hipErrorInvalidValue);
}
SECTION("Pass nullptr to node type") {
hipGraphNode_t memcpyNode;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr , 0));
REQUIRE(hipGraphNodeGetType(memcpyNode, nullptr) == hipErrorInvalidValue);
}
SECTION("Pass invalid node") {
hipGraphNode_t Node = {};
hipGraphNodeType nodeType;
REQUIRE(hipGraphNodeGetType(Node, &nodeType) == hipErrorInvalidValue);
}
}
TEST_CASE("Unit_hipGraphNodeGetType_Functional") {
constexpr size_t N = 1024;
hipGraphNodeType nodeType;
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipEvent_t event;
hipGraphNode_t waiteventNode;
HIP_CHECK(hipEventCreate(&event));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Delete Node and Assign different Node Type") {
HIP_CHECK(hipStreamWaitEvent(stream, event, 0));
HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr , 0));
HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeEmpty);
}
SECTION("Override the graph node and get Node Type") {
HIP_CHECK(hipStreamWaitEvent(stream, event, 0));
HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
HIP_CHECK(hipGraphAddEmptyNode(&waiteventNode, graph, nullptr , 0));
HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeEmpty);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipEventDestroy(event));
}
/**
* Functional Test for hipGraphNodeGetType API
*/
TEST_CASE("Unit_hipGraphNodeGetType_NodeType") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
hipGraphNodeType nodeType;
hipGraphNode_t memcpyNode, kernelNode;
SECTION("Get Memcpy NodeType") {
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphNodeGetType(memcpyNode, &nodeType));
#if HT_AMD
REQUIRE(nodeType == hipGraphNodeTypeMemcpy1D);
#else
REQUIRE(nodeType == cudaGraphNodeTypeMemcpy);
#endif
HIP_CHECK(hipGraphAddEmptyNode(&memcpyNode, graph, nullptr , 0));
HIP_CHECK(hipGraphNodeGetType(memcpyNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeEmpty);
}
SECTION("Get Kernel NodeType") {
hipKernelNodeParams kernelNodeParams{};
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, nullptr,
0, &kernelNodeParams));
HIP_CHECK(hipGraphNodeGetType(kernelNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeKernel);
}
SECTION("Get Empty NodeType") {
hipGraphNode_t emptyNode;
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr , 0));
HIP_CHECK(hipGraphNodeGetType(emptyNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeEmpty);
}
SECTION("Get Memset NodeType") {
hipGraphNode_t memsetNode;
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 10;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphNodeGetType(memsetNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeMemset);
}
SECTION("Get WaitEvent NodeType") {
hipEvent_t event;
hipGraphNode_t waiteventNode;
HIP_CHECK(hipEventCreate(&event));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamWaitEvent(stream, event, 0));
HIP_CHECK(hipGraphAddEventWaitNode(&waiteventNode, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphNodeGetType(waiteventNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeWaitEvent);
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipEventDestroy(event));
}
SECTION("Get EventRecord NodeType") {
hipEvent_t event;
hipGraphNode_t recordeventNode;
HIP_CHECK(hipEventCreate(&event));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipEventRecord(event, stream));
HIP_CHECK(hipGraphAddEventRecordNode(&recordeventNode, graph, nullptr, 0,
event));
HIP_CHECK(hipGraphNodeGetType(recordeventNode, &nodeType));
REQUIRE(nodeType == hipGraphNodeTypeEventRecord);
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipEventDestroy(event));
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
}
tests/catch/unit/texture/CMakeLists.txt
+6
Féach ar an gComhad
@@ -25,6 +25,12 @@ set(TEST_SRC
hipCreateTextureObject_Pitch2D.cc
hipCreateTextureObject_Array.cc
hipTextureObjFetchVector.cc
hipNormalizedFloatValueTex.cc
hipTextureObj2D.cc
hipSimpleTexture3D.cc
hipTextureRef2D.cc
hipSimpleTexture2DLayered.cc
hipTextureMipmapObj2D.cc
)
hip_add_exe_to_target(NAME TextureTest
tests/catch/unit/texture/hipNormalizedFloatValueTex.cc
+166
Féach ar an gComhad
@@ -0,0 +1,166 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#define SIZE 10
#define EPSILON 0.00001
#define THRESH_HOLD 0.01 // For filter mode
static float getNormalizedValue(const float value,
const hipChannelFormatDesc& desc) {
if ((desc.x == 8) && (desc.f == hipChannelFormatKindSigned))
return (value / SCHAR_MAX);
if ((desc.x == 8) && (desc.f == hipChannelFormatKindUnsigned))
return (value / UCHAR_MAX);
if ((desc.x == 16) && (desc.f == hipChannelFormatKindSigned))
return (value / SHRT_MAX);
if ((desc.x == 16) && (desc.f == hipChannelFormatKindUnsigned))
return (value / USHRT_MAX);
return value;
}
texture<char, hipTextureType1D, hipReadModeNormalizedFloat> texc;
texture<unsigned char, hipTextureType1D, hipReadModeNormalizedFloat> texuc;
template<typename T>
__global__ void normalizedValTextureTest(unsigned int numElements,
float* pDst) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
unsigned int elementID = hipThreadIdx_x;
if (elementID >= numElements)
return;
float coord = elementID/static_cast<float>(numElements);
if (std::is_same<T, char>::value)
pDst[elementID] = tex1D(texc, coord);
else if (std::is_same<T, unsigned char>::value)
pDst[elementID] = tex1D(texuc, coord);
#endif
}
static void textureVerifyFilterModePoint(float *hOutputData,
float *expected, int size) {
for (int i = 0; i < size; i++) {
if ((hOutputData[i] == expected[i])
|| (i >= 1 && hOutputData[i] == expected[i - 1]) || // round down
(i < (size - 1) && hOutputData[i] == expected[i + 1])) { // round up
continue;
}
INFO("Mismatch at output[" << i << "]:" << hOutputData[i] <<
" expected[" << i << "]:" << expected[i]);
if (i >= 1) {
INFO(", expected[" << i - 1 << "]:" << expected[i - 1]);
}
if (i < (size - 1)) {
INFO(", expected[" << i + 1 << "]:" << expected[i + 1]);
}
REQUIRE(false);
}
}
static void textureVerifyFilterModeLinear(float *hOutputData,
float *expected, int size) {
for (int i = 0; i < size; i++) {
float mean = (fabs(expected[i]) + fabs(hOutputData[i])) / 2;
float ratio = fabs(expected[i] - hOutputData[i]) / (mean + EPSILON);
if (ratio > THRESH_HOLD) {
INFO("Mismatch found at output[" << i << "]:" << hOutputData[i] <<
" expected[" << i << "]:" << expected[i] << ", ratio:" << ratio);
REQUIRE(false);
}
}
}
template<hipTextureFilterMode fMode = hipFilterModePoint>
static void textureVerify(float *hOutputData, float *expected, size_t size) {
if (fMode == hipFilterModePoint) {
textureVerifyFilterModePoint(hOutputData, expected, size);
} else if (fMode == hipFilterModeLinear) {
textureVerifyFilterModeLinear(hOutputData, expected, size);
}
}
template<typename T, hipTextureFilterMode fMode = hipFilterModePoint>
static void textureTest(texture<T, hipTextureType1D,
hipReadModeNormalizedFloat> *tex) {
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
hipArray_t dData;
HIP_CHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault));
T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74};
HIP_CHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T) * SIZE,
sizeof(T) * SIZE, 1, hipMemcpyHostToDevice));
tex->normalized = true;
tex->channelDesc = desc;
tex->filterMode = fMode;
HIP_CHECK(hipBindTextureToArray(tex, dData, &desc));
float *dOutputData = NULL;
HIP_CHECK(hipMalloc(&dOutputData, sizeof(float) * SIZE));
REQUIRE(dOutputData != nullptr);
hipLaunchKernelGGL(normalizedValTextureTest<T>, dim3(1, 1, 1),
dim3(SIZE, 1, 1), 0, 0, SIZE, dOutputData);
float *hOutputData = new float[SIZE];
REQUIRE(hOutputData != nullptr);
HIP_CHECK(hipMemcpy(hOutputData, dOutputData, (sizeof(float) * SIZE),
hipMemcpyDeviceToHost));
float expected[SIZE];
for (int i = 0; i < SIZE; i++) {
expected[i] = getNormalizedValue(static_cast<float>(hData[i]), desc);
}
textureVerify<fMode>(hOutputData, expected, SIZE);
HIP_CHECK(hipFreeArray(dData));
HIP_CHECK(hipFree(dOutputData));
delete [] hOutputData;
}
template<hipTextureFilterMode fMode = hipFilterModePoint>
static void runTest_hipTextureFilterMode() {
textureTest<char, fMode>(&texc);
textureTest<unsigned char, fMode>(&texuc);
}
TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
hipDeviceProp_t props;
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipGetDeviceProperties(&props, 0));
INFO("Device :: " << props.name);
INFO("Arch - AMD GPU :: " << props.gcnArch);
#endif
SECTION("hipNormalizedFloatValueTexture for hipFilterModePoint") {
runTest_hipTextureFilterMode<hipFilterModePoint>();
}
SECTION("hipNormalizedFloatValueTexture for hipFilterModeLinear") {
runTest_hipTextureFilterMode<hipFilterModeLinear>();
}
}
tests/catch/unit/texture/hipSimpleTexture2DLayered.cc
+103
Féach ar an gComhad
@@ -0,0 +1,103 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
typedef float T;
// Texture reference for 2D Layered texture
texture<float, hipTextureType2DLayered> tex2DL;
__global__ void simpleKernelLayeredArray(T* outputData,
int width, int height, int layer) {
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[layer * width * height + y * width + x] = tex2DLayered(tex2DL,
x, y, layer);
}
TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") {
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
constexpr unsigned int size = width * height * num_layers * sizeof(T);
T* hData = reinterpret_cast<T*>(malloc(size));
REQUIRE(hData != nullptr);
memset(hData, 0, size);
for (unsigned int layer = 0; layer < num_layers; layer++) {
for (int i = 0; i < static_cast<int>(width * height); i++) {
hData[layer * width * height + i] = i;
}
}
hipChannelFormatDesc channelDesc;
// Allocate array and copy image data
channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0,
hipChannelFormatKindFloat);
hipArray *arr;
HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc,
make_hipExtent(width, height, num_layers), hipArrayLayered));
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width , height, num_layers);
// myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
// set texture parameters
tex2DL.addressMode[0] = hipAddressModeWrap;
tex2DL.addressMode[1] = hipAddressModeWrap;
tex2DL.filterMode = hipFilterModePoint;
tex2DL.normalized = false;
// Bind the array to the texture
HIP_CHECK(hipBindTextureToArray(tex2DL, arr, channelDesc));
// Allocate device memory for result
T* dData = nullptr;
HIP_CHECK(hipMalloc(&dData, size));
REQUIRE(dData != nullptr);
dim3 dimBlock(8, 8, 1);
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
for (unsigned int layer = 0; layer < num_layers; layer++)
hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0,
dData, width, height, layer);
HIP_CHECK(hipDeviceSynchronize());
// Allocate mem for the result on host side
T *hOutputData = reinterpret_cast<T*>(malloc(size));
REQUIRE(hOutputData != nullptr);
memset(hOutputData, 0, size);
// copy result from device to host
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
HipTest::checkArray(hData, hOutputData, width, height, num_layers);
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(arr));
free(hData);
free(hOutputData);
}
tests/catch/unit/texture/hipSimpleTexture3D.cc
+127
Féach ar an gComhad
@@ -0,0 +1,127 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
// Texture reference for 3D texture
texture<float, hipTextureType3D, hipReadModeElementType> texf;
texture<int, hipTextureType3D, hipReadModeElementType> texi;
texture<char, hipTextureType3D, hipReadModeElementType> texc;
template <typename T>
__global__ void simpleKernel3DArray(T* outputData, int width,
int height, int depth) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
if (std::is_same<T, float>::value)
outputData[i*width*height + j*width + k] = tex3D(texf, k, j, i);
else if (std::is_same<T, int>::value)
outputData[i*width*height + j*width + k] = tex3D(texi, k, j, i);
else if (std::is_same<T, char>::value)
outputData[i*width*height + j*width + k] = tex3D(texc, k, j, i);
}
}
}
#endif
}
template <typename T>
static void runSimpleTexture3D_Check(int width, int height, int depth,
texture<T, hipTextureType3D, hipReadModeElementType> *tex) {
unsigned int size = width * height * depth * sizeof(T);
T* hData = reinterpret_cast<T*>(malloc(size));
REQUIRE(hData != nullptr);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
hData[i*width*height + j*width +k] = i*width*height + j*width + k;
}
}
}
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray *arr;
HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc,
make_hipExtent(width, height, depth), hipArrayDefault));
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width, height, depth);
myparms.kind = hipMemcpyHostToDevice;
HIP_CHECK(hipMemcpy3D(&myparms));
// set texture parameters
tex->addressMode[0] = hipAddressModeWrap;
tex->addressMode[1] = hipAddressModeWrap;
tex->filterMode = hipFilterModePoint;
tex->normalized = false;
// Bind the array to the texture
HIP_CHECK(hipBindTextureToArray(*tex, arr, channelDesc));
// Allocate device memory for result
T* dData = nullptr;
HIP_CHECK(hipMalloc(&dData, size));
REQUIRE(dData != nullptr);
hipLaunchKernelGGL(simpleKernel3DArray, dim3(1, 1, 1), dim3(1, 1, 1),
0, 0, dData, width, height, depth);
HIP_CHECK(hipDeviceSynchronize());
// Allocate mem for the result on host side
T *hOutputData = reinterpret_cast<T*>(malloc(size));
REQUIRE(hOutputData != nullptr);
memset(hOutputData, 0, size);
// copy result from device to host
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
HipTest::checkArray(hData, hOutputData, width, height, depth);
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(arr));
free(hData);
free(hOutputData);
}
TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
for ( int i = 1; i < 25; i++ ) {
runSimpleTexture3D_Check<float>(i, i, i, &texf);
runSimpleTexture3D_Check<int>(i+1, i, i, &texi);
runSimpleTexture3D_Check<char>(i, i+1, i, &texc);
}
}
tests/catch/unit/texture/hipTextureMipmapObj2D.cc
+153
Féach ar an gComhad
@@ -0,0 +1,153 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
// Height Width Vector
std::vector<unsigned int> hw_vector = {2048, 1024, 512, 256, 64};
std::vector<unsigned int> mip_vector = {8, 4, 2, 1};
__global__ void tex2DKernel(float* outputData,
hipTextureObject_t textureObject,
int width, float level) {
#ifndef __gfx90a__
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[y * width + x] = tex2DLod<float>(textureObject, x, y, level);
#endif
#endif
}
#ifdef _WIN32 // MipMap is currently supported only on windows
static void runMipMapTest(unsigned int width, unsigned int height,
unsigned int mipmap_level) {
INFO("Width: " << width << "Height: " << height << "mip: " << mipmap_level);
// Create new width & height to be tested
unsigned int orig_width = width;
unsigned int orig_height = height;
unsigned int i, j;
width /= pow(2, mipmap_level);
height /= pow(2, mipmap_level);
unsigned int size = width * height * sizeof(float);
float* hData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hData != nullptr);
memset(hData, 0, size);
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0,
hipChannelFormatKindFloat);
HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc;
memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR));
mipmapped_array_desc.Width = orig_width;
mipmapped_array_desc.Height = orig_height;
mipmapped_array_desc.Depth = 0;
mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT;
mipmapped_array_desc.NumChannels = ((channelDesc.x != 0) +
(channelDesc.y != 0) + (channelDesc.z != 0) + (channelDesc.w != 0));
mipmapped_array_desc.Flags = 0;
hipMipmappedArray* mip_array_ptr;
HIP_CHECK(hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc,
2 * mipmap_level));
hipArray *hipArray = nullptr;
HIP_CHECK(hipMipmappedArrayGetLevel(&hipArray, mip_array_ptr, mipmap_level));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Specify texture object parameters
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = hipAddressModeWrap;
texDesc.addressMode[1] = hipAddressModeWrap;
texDesc.filterMode = hipFilterModePoint;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = 0;
// Create texture object
hipTextureObject_t textureObject = 0;
HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc,
&texDesc, nullptr));
float* dData = nullptr;
HIP_CHECK(hipMalloc(&dData, size));
REQUIRE(dData != nullptr);
dim3 dimBlock(16, 16, 1);
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData,
textureObject, width, (2 * mipmap_level));
hipDeviceSynchronize();
float* hOutputData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hOutputData != nullptr);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
INFO("Difference found at [ " << i << j << " ]: " <<
hData[i * width + j] << hOutputData[i * width + j]);
REQUIRE(false);
}
}
}
HIP_CHECK(hipDestroyTextureObject(textureObject));
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(hipArray));
free(hData);
}
#endif
TEST_CASE("Unit_hipTextureMipmapObj2D_Check") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
#ifdef _WIN32
for (auto& hw : hw_vector) {
for (auto& mip : mip_vector) {
if ((hw / static_cast<int>(pow(2, (mip * 2)))) > 0) {
runMipMapTest(hw, hw, mip);
}
}
}
#else
SUCCEED("Mipmaps are Supported only on windows, skipping the test.");
#endif
}
tests/catch/unit/texture/hipTextureObj2D.cc
+113
Féach ar an gComhad
@@ -0,0 +1,113 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
__global__ void tex2DKernel(float* outputData,
hipTextureObject_t textureObject, int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[y * width + x] = tex2D<float>(textureObject, x, y);
#endif
}
TEST_CASE("Unit_hipTextureObj2D_Check") {
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
constexpr unsigned int size = width * height * sizeof(float);
unsigned int i, j;
float* dData = nullptr;
HIP_CHECK(hipMalloc(&dData, size));
REQUIRE(dData != nullptr);
float* hOutputData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hOutputData != nullptr);
memset(hOutputData, 0, size);
float* hData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hData != nullptr);
memset(hData, 0, size);
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
}
hipChannelFormatDesc channelDesc =
hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
hipArray* hipArray;
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = hipArray;
// Specify texture object parameters
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = hipAddressModeWrap;
texDesc.addressMode[1] = hipAddressModeWrap;
texDesc.filterMode = hipFilterModePoint;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = 0;
// Create texture object
hipTextureObject_t textureObject = 0;
HIP_CHECK(hipCreateTextureObject(&textureObject, &resDesc,
&texDesc, nullptr));
dim3 dimBlock(16, 16, 1);
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock),
0, 0, dData, textureObject, width);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
INFO("Difference found at [ " << i << j << " ]: " <<
hData[i * width + j] << hOutputData[i * width + j]);
REQUIRE(false);
}
}
}
HIP_CHECK(hipDestroyTextureObject(textureObject));
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(hipArray));
free(hData);
}
tests/catch/unit/texture/hipTextureRef2D.cc
+87
Féach ar an gComhad
@@ -0,0 +1,87 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
texture<float, 2, hipReadModeElementType> tex;
__global__ void tex2DKernel(float* outputData, int width) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[y * width + x] = tex2D(tex, x, y);
}
TEST_CASE("Unit_hipTextureRef2D_Check") {
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
constexpr unsigned int size = width * height * sizeof(float);
unsigned int i, j;
float* hData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hData != nullptr);
memset(hData, 0, size);
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
}
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0,
hipChannelFormatKindFloat);
hipArray* hipArray;
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, height));
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size,
hipMemcpyHostToDevice));
tex.addressMode[0] = hipAddressModeWrap;
tex.addressMode[1] = hipAddressModeWrap;
tex.filterMode = hipFilterModePoint;
tex.normalized = 0;
HIP_CHECK(hipBindTextureToArray(tex, hipArray, channelDesc));
float* dData = nullptr;
HIP_CHECK(hipMalloc(&dData, size));
REQUIRE(dData != nullptr);
dim3 dimBlock(16, 16, 1);
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
dData, width);
hipDeviceSynchronize();
float* hOutputData = reinterpret_cast<float*>(malloc(size));
REQUIRE(hOutputData != nullptr);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
for (i = 0; i < height; i++) {
for (j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
INFO("Difference found at [ " << i << j << " ]: " <<
hData[i * width + j] << hOutputData[i * width + j]);
REQUIRE(false);
}
}
}
HIP_CHECK(hipUnbindTexture(tex));
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(hipArray));
}
tests/hit/HIT.cmake
+13 -11
Féach ar an gComhad
@@ -45,19 +45,21 @@ if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerat
execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH
RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT)
message(STATUS "ROCm Agent Enumurator Result: ${ROCM_AGENT_ENUM_RESULT}")
# Trim out gfx000
string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH})
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
if (NOT HIP_GPU_ARCH STREQUAL "")
string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN)
# If string has more gfx target except gfx000
if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1)
string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}")
set(OFFLOAD_ARCH_STR "")
foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
endforeach()
message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
endif()
else()
message(STATUS "ROCm Agent Enumurator found no valid architectures")
endif()
else()
message(STATUS "ROCm Agent Enumurator Not Found")