AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

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

Kaynağa Gözat 
Change-Id: Idef3b509e0aeaf57e2a58906acd2538d3106a1c3
Bu işleme şunda yer alıyor:
Jenkins
2022-06-20 13:03:21 -04:00
ebeveyn 1dfaa6d463 cc96fdec02
işleme 80caf65b19
82 değiştirilmiş dosya ile 4027 ekleme ve 444 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
.jenkins/Jenkinsfile
Sağlanmış
+1 -1
Dosyayı Görüntüle
@@ -101,7 +101,7 @@ def hipBuildTest(String backendLabel) {
# Check if backend label contains string "amd" or backend host is a server with amd gpu
if [[ $backendLabel =~ amd ]]; then
export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json"
LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple'
LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2'
sleep 120
else
make test
bin/hipcc
+7 -1
Dosyayı Görüntüle
@@ -22,7 +22,6 @@
# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;
use strict;
use warnings;
use File::Basename;
@@ -30,6 +29,13 @@ use File::Spec::Functions 'catfile';
# TODO: By default select perl script until change incorporated in HIP build script.
my $HIPCC_USE_PERL_SCRIPT = 1;
my $isWindows = ($^O eq 'MSWin32' or $^O eq 'msys');
# escapes args with quotes SWDEV-341955
foreach $arg (@ARGV) {
if ($isWindows) {
$arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g;
}
}
my $SCRIPT_DIR=dirname(__FILE__);
if ($HIPCC_USE_PERL_SCRIPT) {
bin/hipcc.pl
+4 -1
Dosyayı Görüntüle
@@ -595,9 +595,12 @@ foreach $arg (@ARGV)
# common characters such as alphanumerics.
# Do the quoting here because sometimes the $arg is changed in the loop
# Important to have all of '-Xlinker' in the set of unquoted characters.
if (not $isWindows and $escapeArg) { # Windows needs different quoting, ignore for now
if (not $isWindows and $escapeArg) {
$arg =~ s/[^-a-zA-Z0-9_=+,.\/]/\\$&/g;
}
if ($isWindows and $escapeArg) {
$arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g;
}
$toolArgs .= " $arg" unless $swallowArg;
$prevArg = $arg;
}
docs/markdown/hip_build.md
+77 -5
Dosyayı Görüntüle
@@ -8,10 +8,12 @@
* [Set the environment variables](#set-the-environment-variables)
* [Build HIP](#build-hip)
* [Default paths and environment variables](#default-paths-and-environment-variables)
* [Build HIP Tests](#build-hip-tests)
- [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)
* [Build HIP tests](#build-hip-tests)
- [Run HIP](#run-hip)
<!-- tocstop -->
@@ -30,7 +32,6 @@ 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)
@@ -61,7 +62,6 @@ ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is at /opt/rocm.
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
@@ -69,8 +69,6 @@ git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtim
```
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.
@@ -84,7 +82,7 @@ See https://github.com/ROCm-Developer-Tools/hipamd
```
cd "$HIPAMD_DIR"
mkdir -p build; cd build
cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=$OPENCL_DIR -DROCCLR_PATH=$ROCCLR_DIR -DCMAKE_PREFIX_PATH="<ROCM_PATH>/" -DCMAKE_INSTALL_PREFIX=$PWD/install ..
cmake -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_PREFIX_PATH="<ROCM_PATH>/" -DCMAKE_INSTALL_PREFIX=$PWD/install ..
make -j$(nproc)
sudo make install
```
@@ -103,6 +101,77 @@ By default, release version of AMDHIP is built.
After make install command, make sure HIP_PATH is pointed to $PWD/install/hip.
## Build HIP tests
### Build HIP directed tests
Developers can build HIP directed tests right after build HIP commands,
```
sudo make install
make -j$(nproc) build_tests
```
By default, all HIP directed tests will be built and generated under the folder $HIPAMD_DIR/build/directed_tests.
Take HIP directed device APIs tests, as an example, all available test applications will have executable files generated under,
$HIPAMD_DIR/build/directed_tests/runtimeApi/device.
Run all HIP directed_tests, use the command,
```
ctest
```
Or
```
make test
```
Build and run a single directed test, use the follow command as an example,
```
make directed_tests.texture.hipTexObjPitch
cd $HIPAMD_DIR/build/directed_tests/texcture
./hipTexObjPitch
```
Please note, the integrated HIP directed tests, will be deprecated in future release.
### Build HIP catch tests
After build and install HIP commands, catch tests can be built via the following instructions,
```
cd "$HIP_DIR"
mkdir -p build; cd build
export HIP_PATH=$HIPAMD_DIR/build
cmake ../tests/catch/ -DHIP_PLATFORM=amd
make -j$(nproc) build_tests
ctest # run tests
```
HIP catch tests are built under the folder $HIP_DIR/build.
To run a single catch test, the following is an example,
```
cd $HIP_DIR/build/unit/texture
./TextureTest
```
### Build HIP Catch2 standalone test
HIP Catch2 supports build a standalone test, for example,
```
export PATH=$HIP_DIR/bin:$PATH
export HIP_PATH=$HIPAMD_DIR/build/install
hipcc $HIP_DIR/tests/catch/unit/memory/hipPointerGetAttributes.cc -I ./tests/catch/include ./tests/catch/hipTestMain/standalone_main.cc -I ./tests/catch/external/Catch2 -g -o hipPointerGetAttributes
./hipPointerGetAttributes
...
All tests passed
```
HIP catch tests, especially new architectured Catch2, will be official HIP tests in the repository and can be built alone as with the instructions shown above.
# Build HIP on NVIDIA platform
@@ -131,6 +200,9 @@ make -j$(nproc)
sudo make install
```
## Build HIP tests
Build HIP tests commands on NVIDIA platform are basically the same as AMD, except set -DHIP_PLATFORM=nvidia.
# Run HIP
Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-5.0.x/samples/0_Intro/square).
hip-lang-config.cmake.in
+2 -2
Dosyayı Görüntüle
@@ -94,7 +94,7 @@ find_path(HSA_HEADER hsa/hsa.h
/opt/rocm/include
)
if (HSA_HEADER-NOTFOUND)
if (NOT HSA_HEADER)
message (FATAL_ERROR "HSA header not found! ROCM_PATH environment not set")
endif()
@@ -136,7 +136,7 @@ set_property(TARGET hip-lang::device APPEND PROPERTY
)
# Add support for __fp16 and _Float16, explicitly link with compiler-rt
if(CLANGRT_BUILTINS-NOTFOUND)
if(NOT CLANGRT_BUILTINS)
message(FATAL_ERROR "clangrt builtins lib not found")
else()
set_property(TARGET hip-lang::device APPEND PROPERTY
include/hip/hip_runtime_api.h
+4 -3
Dosyayı Görüntüle
@@ -1192,7 +1192,6 @@ typedef enum hipGraphInstantiateFlags {
hipGraphInstantiateFlagAutoFreeOnLaunch =
1, ///< Automatically free memory allocated in a graph before relaunching.
} hipGraphInstantiateFlags;
#include <hip/amd_detail/amd_hip_runtime_pt_api.h>
// Doxygen end group GlobalDefs
/** @} */
@@ -1517,9 +1516,9 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
*/
hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig);
/**
* @brief Set Cache configuration for a specific function
* @brief Get Cache configuration for a specific Device
*
* @param [in] cacheConfig
* @param [out] cacheConfig
*
* @returns #hipSuccess, #hipErrorNotInitialized
* Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored
@@ -6813,6 +6812,8 @@ static inline hipError_t hipMallocManaged(T** devPtr, size_t size,
#endif
#endif
#include <hip/amd_detail/amd_hip_runtime_pt_api.h>
#if USE_PROF_API
#include <hip/amd_detail/hip_prof_str.h>
#endif
include/hip/hiprtc.h
+23
Dosyayı Görüntüle
@@ -269,6 +269,29 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
*/
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
/**
* @brief Gets the pointer of compiled bitcode by the runtime compilation program instance.
*
* @param [in] prog runtime compilation program instance.
* @param [out] code char pointer to bitcode.
* @return HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode);
/**
* @brief Gets the size of compiled bitcode by the runtime compilation program instance.
*
*
* @param [in] prog runtime compilation program instance.
* @param [out] code the size of bitcode.
* @return HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size);
/**
* @brief Creates the link instance via hiprtc APIs.
*
samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp
+2 -2
Dosyayı Görüntüle
@@ -37,8 +37,8 @@ THE SOFTWARE.
// Device (Kernel) function, it must be void
__global__ void matrixTranspose(float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
out[y * width + x] = in[x * width + y];
}
samples/2_Cookbook/1_hipEvent/hipEvent.cpp
+2 -2
Dosyayı Görüntüle
@@ -35,8 +35,8 @@ THE SOFTWARE.
// Device (Kernel) function, it must be void
__global__ void matrixTranspose(float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
out[y * width + x] = in[x * width + y];
}
samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt
+97
Dosyayı Görüntüle
@@ -0,0 +1,97 @@
# 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.
# hipcc.bat fails to qualify as a valid compiler for CMAKE_CXX_COMPILER_ID = ROCMClang
# so the simple compiler test is skipped and forced to use hipcc.bat as compiler
set(CMAKE_C_COMPILER_WORKS 1)
set(CMAKE_CXX_COMPILER_WORKS 1)
set(CMAKE_CXX_STANDARD 14)
project(hipVulkan)
cmake_minimum_required(VERSION 3.10)
set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake;${CMAKE_MODULE_PATH}")
if (NOT DEFINED ROCM_PATH )
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
endif ()
# Search for rocm in common locations
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
# need to set rocm_path for windows
# since clang and hip are two different folders during build/install step
if (WIN32 AND HIPINFO_INTERNAL_BUILD)
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${HIP_PATH}")
endif()
# Find hip
find_package(hip REQUIRED)
if (WIN32)
find_package(GLFW3)
if(NOT GLFW_FOUND)
if(EXISTS "${GLFW_PATH}")
message(STATUS "FOUND GLFW SDK: ${GLFW_PATH}")
elseif (EXISTS "$ENV{GLFW_PATH}")
message(STATUS "FOUND GLFW SDK: $ENV{GLFW_PATH}")
set(GLFW_PATH $ENV{GLFW_PATH})
else()
message("Error: Unable to locate GLFW SDK. please specify GLFW_PATH")
return()
endif()
endif()
endif(WIN32)
find_package(Vulkan)
if(NOT Vulkan_FOUND)
if(EXISTS "${VULKAN_PATH}")
message(STATUS "Vulkan SDK: ${VULKAN_PATH}")
elseif (EXISTS "$ENV{VULKAN_SDK}")
message(STATUS "FOUND VULKAN SDK: $ENV{VULKAN_SDK}")
set(VULKAN_PATH $ENV{VULKAN_SDK})
else()
message("Error: Unable to locate Vulkan SDK. please specify VULKAN_PATH")
return()
endif()
endif()
set(VULKAN_PATH ${Vulkan_INCLUDE_DIRS})
STRING(REGEX REPLACE "/[Ii]nclude" "" VULKAN_PATH ${VULKAN_PATH})
# Include Vulkan header files from Vulkan SDK
include_directories(AFTER ${VULKAN_PATH}/include)
link_directories(${VULKAN_PATH}/bin;${VULKAN_PATH}/lib;)
link_directories(${GLFW_PATH}/lib-vc2019)
# Set compiler and linker
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
set(CMAKE_BUILD_TYPE Release)
# Create the excutable
add_executable(hipVulkan VulkanBaseApp.cpp VulkanBaseApp.h main.cpp SineWaveSimulation.cpp SineWaveSimulation.h linmath.h)
include_directories(${HIP_PATH}/include)
include_directories(${GLFW_PATH}/include)
# Link with HIP
if (WIN32)
target_link_libraries(hipVulkan advapi32 hip::host vulkan-1 glfw3dll)
else (WIN32)
target_link_libraries(hipVulkan hip::host vulkan glfw)
endif (WIN32)
samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp
+147
Dosyayı Görüntüle
@@ -0,0 +1,147 @@
/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
* * Neither the name of NVIDIA CORPORATION nor the names of its
* contributors may be used to endorse or promote products derived
* from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/
/*
* Modifications Copyright (C)2021 Advanced
* Micro Devices, Inc. All rights reserved.
*/
#include "SineWaveSimulation.h"
#include <algorithm>
//#include <helper_cuda.h>
#include "hip/hip_runtime.h"
__global__ void sinewave(float *heightMap, unsigned int width, unsigned int height, float time)
{
const float freq = 4.0f;
const size_t stride = gridDim.x * blockDim.x;
// Iterate through the entire array in a way that is
// independent of the grid configuration
for (size_t tid = blockIdx.x * blockDim.x + threadIdx.x; tid < width * height; tid += stride) {
// Calculate the x, y coordinates
const size_t y = tid / width;
const size_t x = tid - y * width;
// Normalize x, y to [0,1]
const float u = ((2.0f * x) / width) - 1.0f;
const float v = ((2.0f * y) / height) - 1.0f;
// Calculate the new height value
const float w = 0.5f * sinf(u * freq + time) * cosf(v * freq + time);
// Store this new height value
heightMap[tid] = w;
}
}
SineWaveSimulation::SineWaveSimulation(size_t width, size_t height)
: m_heightMap(nullptr), m_width(width), m_height(height)
{
}
void SineWaveSimulation::initCudaLaunchConfig(int device)
{
hipDeviceProp_t prop = {};
checkHIPErrors(hipSetDevice(device));
checkHIPErrors(hipGetDeviceProperties(&prop, device));
// We don't need large block sizes, since there's not much inter-thread communication
m_threads = prop.warpSize;
// Use the occupancy calculator and fill the gpu as best as we can
checkHIPErrors(hipOccupancyMaxActiveBlocksPerMultiprocessor(&m_blocks, sinewave, prop.warpSize, 0));
m_blocks *= prop.multiProcessorCount;
// Go ahead and the clamp the blocks to the minimum needed for this height/width
m_blocks = std::min(m_blocks, (int)((m_width * m_height + m_threads - 1) / m_threads));
}
int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE)
{
int current_device = 0;
int device_count = 0;
int devices_prohibited = 0;
hipDeviceProp_t deviceProp;
checkHIPErrors(hipGetDeviceCount(&device_count));
if (device_count == 0) {
fprintf(stderr, "CUDA error: no devices supporting CUDA.\n");
exit(EXIT_FAILURE);
}
// Find the GPU which is selected by Vulkan
while (current_device < device_count) {
hipGetDeviceProperties(&deviceProp, current_device);
if ((deviceProp.computeMode != hipComputeModeProhibited)) {
// Compare the cuda device UUID with vulkan UUID
// FIXME
int ret = 0; // memcmp((void*)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE);
if (ret == 0)
{
checkHIPErrors(hipSetDevice(current_device));
checkHIPErrors(hipGetDeviceProperties(&deviceProp, current_device));
printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n",
current_device, deviceProp.name, deviceProp.major,
deviceProp.minor);
return current_device;
}
} else {
devices_prohibited++;
}
current_device++;
}
if (devices_prohibited == device_count) {
fprintf(stderr,
"HIP error:"
" No Vulkan-HIP Interop capable GPU found.\n");
exit(EXIT_FAILURE);
}
return -1;
}
SineWaveSimulation::~SineWaveSimulation()
{
m_heightMap = NULL;
}
void SineWaveSimulation::initSimulation(float *heights)
{
m_heightMap = heights;
}
void SineWaveSimulation::stepSimulation(float time, hipStream_t stream)
{
hipLaunchKernelGGL(sinewave, dim3(m_blocks), dim3(m_threads), 0, stream , m_heightMap, m_width, m_height, time);
getLastHIPError("Failed to launch CUDA simulation");
//hipStreamSynchronize(stream);
}
samples/2_Cookbook/20_hip_vulkan/buildcmd.txt
+6 -1
Dosyayı Görüntüle
@@ -5,7 +5,12 @@
o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv
o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv
to build without cmake:
• set HCC_AMDGPU_TARGET=gfx906:sramecc-:xnack- (for your graphic card, you can get the name from hipinfo )
$• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14
• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14
• run a.exe, you should see a 3D sinewave simulation
to build with cmake on windows:
• mkdir build; cd build
• cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=d:\driver2\drivers\drivers\compute\hip_sdk
samples/2_Cookbook/20_hip_vulkan/main.cpp
+14 -2
Dosyayı Görüntüle
@@ -37,7 +37,7 @@
#include <chrono>
#include <algorithm>
#include "linmath.h"
#include "hip_runtime.h"
#include "hip/hip_runtime.h"
#include "SineWaveSimulation.h"
@@ -50,6 +50,18 @@ std::string execution_path;
#define ENABLE_VALIDATION (true)
#endif
#ifndef _WIN64
#define MAX_PATH 260
int GetModuleFileName(void* hndl, char* name, int size)
{
FILE* stream = fopen("/proc/self/cmdline", "r");
fgets(name, size, stream);
fclose(stream);
return strlen(name);
}
#endif
class VulkanCudaSineWave : public VulkanBaseApp
{
@@ -93,7 +105,7 @@ public:
throw std::runtime_error("Requested height and width is too large for this sample!");
}
// Add our compiled vulkan shader files
TCHAR buffer[MAX_PATH] = { 0 };
char buffer[MAX_PATH] = { 0 }; //assuming none unicode
GetModuleFileName(NULL, buffer, MAX_PATH);
std::string str3 = std::string(buffer);
std::string str1 = "vert.spv" ; //sdkFindFilePath("sinewave.vert", execution_path.c_str());
samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt
+34
Dosyayı Görüntüle
@@ -0,0 +1,34 @@
# 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.
cmake_minimum_required(VERSION 3.21.3)
project(cmake_cxx_amdclang++_test
DESCRIPTION "Verifies CXX Language Support with amdclang++"
LANGUAGES CXX)
# Find HIP
find_package(hip REQUIRED)
# Create the excutable
add_executable(square square.cpp)
# Link with HIP
target_link_libraries(square hip::device)
samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md
+19
Dosyayı Görüntüle
@@ -0,0 +1,19 @@
### This sample tests CXX Language support with amdclang++
I. Build
mkdir -p build; cd build
rm -rf *;
CXX=`hipconfig -l`/amdclang++ cmake .. (or)
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ .. (or)
cmake -DCMAKE_CXX_COMPILER=/opt/rocm-X.Y.Z/llvm/bin/amdclang++ ..
make
II. Test
$ ./square
info: running on device Vega 20 [Radeon Pro Vega 20]
info: allocate host mem ( 7.63 MB)
info: allocate device mem ( 7.63 MB)
info: copy Host2Device
info: launch 'vector_square' kernel
info: copy Device2Host
info: check result
PASSED!
samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp
+98
Dosyayı Görüntüle
@@ -0,0 +1,98 @@
/*
Copyright (c) 2015-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 <stdio.h>
#include <hip/hip_runtime.h>
#define CHECK(cmd) \
{\
hipError_t error = cmd;\
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
exit(EXIT_FAILURE);\
}\
}
/*
* Square each element in the array A and write to array C.
*/
template <typename T>
__global__ void
vector_square(T *C_d, T *A_d, size_t N)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
for (size_t i=offset; i<N; i+=stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
int main(int argc, char *argv[])
{
float *A_d, *C_d;
float *A_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
printf ("info: running on device %s\n", props.name);
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess );
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess );
// Fill with Phi + i
for (size_t i=0; i<N; i++)
{
A_h[i] = 1.618f + i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: copy Host2Device\n");
CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
printf ("info: launch 'vector_square' kernel\n");
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf ("info: check result\n");
for (size_t i=0; i<N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
}
}
printf ("PASSED!\n");
}
samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt
+27
Dosyayı Görüntüle
@@ -0,0 +1,27 @@
# 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.
cmake_minimum_required(VERSION 3.21.3)
project(cmake_hip_lang_support VERSION 1.0
DESCRIPTION "HIP Language Support with upstream CMake"
LANGUAGES HIP)
# Create the executable
add_executable(square square.hip)
samples/2_Cookbook/22_cmake_hip_lang/README.md
+16
Dosyayı Görüntüle
@@ -0,0 +1,16 @@
### This will test HIP language support in upstream CMake
I. Build
mkdir -p build; cd build
rm -rf *; cmake ..
make
II. Test
$ ./square
info: running on device
info: allocate host mem ( 7.63 MB)
info: allocate device mem ( 7.63 MB)
info: copy Host2Device
info: launch 'vector_square' kernel
info: copy Device2Host
info: check result
PASSED!
samples/2_Cookbook/22_cmake_hip_lang/square.hip
+98
Dosyayı Görüntüle
@@ -0,0 +1,98 @@
/*
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 <stdio.h>
#include <hip/hip_runtime.h>
#define CHECK(cmd) \
{\
hipError_t error = cmd;\
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
exit(EXIT_FAILURE);\
}\
}
/*
* Square each element in the array A and write to array C.
*/
template <typename T>
__global__ void
vector_square(T *C_d, T *A_d, size_t N)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
for (size_t i=offset; i<N; i+=stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
int main(int argc, char *argv[])
{
float *A_d, *C_d;
float *A_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/));
printf ("info: running on device %s\n", props.name);
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess );
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess );
// Fill with Phi + i
for (size_t i=0; i<N; i++)
{
A_h[i] = 1.618f + i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: copy Host2Device\n");
CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
printf ("info: launch 'vector_square' kernel\n");
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf ("info: check result\n");
for (size_t i=0; i<N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
CHECK(hipErrorUnknown);
}
}
printf ("PASSED!\n");
}
samples/2_Cookbook/3_shared_memory/sharedMemory.cpp
+2 -2
Dosyayı Görüntüle
@@ -38,8 +38,8 @@ THE SOFTWARE.
__global__ void matrixTranspose(float* out, float* in, const int width) {
__shared__ float sharedMem[WIDTH * WIDTH];
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
sharedMem[y * width + x] = in[x * width + y];
samples/2_Cookbook/4_shfl/shfl.cpp
+1 -1
Dosyayı Görüntüle
@@ -36,7 +36,7 @@ THE SOFTWARE.
// Device (Kernel) function, it must be void
__global__ void matrixTranspose(float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int x = blockDim.x * blockIdx.x + threadIdx.x;
float val = in[x];
samples/2_Cookbook/5_2dshfl/2dshfl.cpp
+2 -2
Dosyayı Görüntüle
@@ -36,8 +36,8 @@ THE SOFTWARE.
// Device (Kernel) function, it must be void
__global__ void matrixTranspose(float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
float val = in[y * width + x];
out[x * width + y] = __shfl(val, y * width + x);
samples/2_Cookbook/5_2dshfl/Readme.md
+1 -1
Dosyayı Görüntüle
@@ -31,7 +31,7 @@ We will be using the Simple Matrix Transpose application from the previous tutor
In the same sourcecode, we used for MatrixTranspose. We'll add the following:
```
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
int y = blockDim.y * blockIdx.y + threadIdx.y;
out[x*width + y] = __shfl(val,y*width + x);
```
tests/README.md
+16
Dosyayı Görüntüle
@@ -147,6 +147,22 @@ For example,
Here "-C performance" indicate the "performance" configuration of ctest.
```
### RTC Testing
To enable RTC testing, cmake needs to be passed the DRTC_TESTING=1 options.
When this option is passed, all tests that support this functionality will be run using HIP RTC to compile and run.
To enable HIP RTC support for a specific test:
1 - Move all its kernels to tests/catch/kernels (one file per kernel)
2 - Update tests/catch/kernels/CMakeLists.txt
3 - Update tests/catch/include/kernels.hh
4 - Update tests/catch/include/kernel_mapping.hh
5 - Include kernels.hh
6 - Call hipTest::launchKernel() function instead of hipLaunchKernelGGL()
Note: HIP RTC does not do implicit casting of kernel parameters. This requires the test writer to explicitly do all the casting before running the kernel. The code will not compile otherwise.
### If a test fails - how to debug a test
Find the test and commandline that fail:
tests/catch/CMakeLists.txt
+8
Dosyayı Görüntüle
@@ -104,10 +104,17 @@ set(CATCH2_INCLUDE ${CATCH2_PATH}/cmake/Catch2/catch_include.cmake.in)
include_directories(
${CATCH2_PATH}
"./include"
"./kernels"
${HIP_PATH}/include
${JSON_PARSER}
)
option(RTC_TESTING "Run tests using HIP RTC to compile the kernels" OFF)
if (RTC_TESTING)
add_definitions(-DRTC_TESTING=ON)
endif()
add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/")
file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/hipTestMain)
file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/script)
set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake)
@@ -177,6 +184,7 @@ add_custom_target(build_tests)
# Tests folder
add_subdirectory(unit)
add_subdirectory(ABM)
add_subdirectory(kernels)
add_subdirectory(hipTestMain)
add_subdirectory(stress)
add_subdirectory(TypeQualifiers)
tests/catch/README.md
+112 -13
Dosyayı Görüntüle
@@ -34,6 +34,22 @@ Some useful functions are:
This information can be accessed in any test via using: `TestContext::get().isAmd()`.
## Adding test for a specific platform
There might be some functionality which is not present on some platforms. Those tests can be hidden inside following macros.
- ```HT_AMD``` is 1 when tests are running on AMD platform and 0 on NVIDIA.
- ```HT_NVIDIA``` is 1 when tests are running on NVIDIA platform and 0 on AMD
Usage:
```cpp
#if HT_AMD
TEST_CASE("hipExtAPIs") {
// ...
}
#endif
```
## Config file schema
Some tests can be skipped using a config file placed in hipTestMain/config folder. Multiple config files can be defined for different configurations.
The naming convention for the file needs to be "config_platform_os_archname.json"
@@ -56,16 +72,108 @@ The schema of the json file is as follows:
}
```
## Env Variables
## Environment Variables
- `HT_CONFIG_FILE` : This variable can be set to the config file name or full path. Disabled tests will be read from this.
- `HT_LOG_ENABLE` : This is for debugging the HIP Test Framework itself. Setting it to 1, all `LogPrintf` will be printed on screen
## Test Macros
### Single Thread Macros
These macros are to be used when your test is calling HIP APIs via the main thread.
- `HIP_CHECK` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```.
- Usage: ```HIP_CHECK(hipMalloc(&dPtr, 10));```
- ```HIP_CHECK_ERROR``` : This macro takes in a HIP API and tests its result against a provided result. This can be used when the API is expected to fail with a particular result.
- Usage: ```HIP_CHECK_ERROR(hipMalloc(&dPtr, 0), hipErrorInvalidValue);```
- ```HIPRTC_CHECK``` : This macro takes in a HIPRTC API and tests its result against HIPRTC_SUCCESS.
- Usage: ```HIPRTC_CHECK(hiprtcCompileProgram(prog, count, options));```
- ```HIP_ASSERT``` : This macro takes in a bool condition as input and does a ```REQUIRE``` on the condition.
- Usage: ```HIP_ASSERT(result == 10);```
### Multi Thread Macros
These macros are to be used when you call HIP APIs in a multi threaded way. They exist because Catch2 ```REQUIRE``` and ```CHECK``` macros can not handle multi threaded calls. To solve this problem, two macros are added```HIP_CHECK_THREAD``` and ```REQUIRE_THREAD``` which can be used to check result of HIP APIs and test assertions respectively. The results can be validate after the threads join via ```HIP_CHECK_THREAD_FINALIZE```.
Note: These should used in ```std::thread``` only. For multi proc guidelines look at [MultiProc Macros](#multi-process-macros) and [SpawnProc Class](#multiproc-management-class)
- ```HIP_CHECK_THREAD``` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. It can also tell other threads if an error has occured in one of the HIP API and can prematurely stop the threads.
- ```REQUIRE_THREAD``` : This macro takes in a bool condition and tests for its result to be true. If this check fails, it can signal other threads to terminate early.
- ```HIP_CHECK_THREAD_FINALIZE``` : This macro checks for the results logged by ```HIP_CHECK_THREAD```. This needs to be called after the threads have joined.
Please also note that you can not return values in functions calling ```HIP_CHECK_THREAD``` or ```REQUIRE_THREAD``` macro.
Usage:
```cpp
auto threadFunc = []() {
int *dPtr{nullptr};
HIP_CHECK_THREAD(hipMalloc(&dPtr, 10));
REQUIRE_THREAD(dPtr != nullptr);
// Some other work
};
// Launch threads
std::vector<std::thread> threadPool;
for(...) {
threadPool.emplace_back(std::thread(threadFunc));
}
// Join threads
for(auto &i : threadPool) {
i.join();
}
// Validate all results
HIP_CHECK_THREAD_FINALIZE();
```
### Skipping Tests if certain criteria is not met
If there arises a condition where certain flag is disabled and due to which a test can not run at that time, the following macro can be of use. It will highlight the test in ctest report as well.
- ```HIP_SKIP_TEST``` : The api takes in an input of the reason as well and prints out the line HIP_SKIP_THIS_TEST. This causes ctest to mark the test as skipped and the test shows up in the report as skipped prompting proper response from the team.
Usage:
```cpp
TEST_CASE("TestOnlyOnXnack") {
if(!XNACKEnabled) {
HIP_SKIP_TEST("Test only runs on system with XNACK enabled");
return;
}
// Rest of test functionality
}
```
### Multi Process Macros
These macros are to be called in multi process tests, inside a process which gets spawned. The reasoning is the same, Catch2 does not support multi process checks.
- ```HIPCHECK``` : Same as ```HIP_CHECK``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process.
- ```HIPASSERT``` : Same as ```HIP_ASSERT``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process.
## MultiProc Management Class
There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output.
Sample Usage:
```cpp
hip::SpawnProc proc(<relative path of exe with test folder>, <optional bool value, if output is to be recorded>);
REQUIRE(0 == proc.run()); // Test of return value of the proc
REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process
```
The process can be a standalone exe (see tests/catch/unit/printfExe for more information).
## Enabling New Tests
Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=ON```. After porting existing tests, this will be turned on by default.
Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default.
## Building a single test
```bash
hipcc <path_to_test.cpp> -I<HIP_SRC_DIR>/tests/newTests/include <HIP_SRC_DIR>/tests/newTests/hipTestMain/standalone_main.cc -I<HIP_SRC_DIR>/tests/newTests/external/Catch2 -g -o <out_file_name>
hipcc <path_to_test.cpp> -I<HIP_SRC_DIR>/tests/catch/include <HIP_SRC_DIR>/tests/catch/hipTestMain/standalone_main.cc -I<HIP_SRC_DIR>/tests/catch/external/Catch2 -g -o <out_file_name>
```
## Debugging support
@@ -87,16 +195,7 @@ Tests fall in 5 categories and its file name prefix are as follows:
- Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment
- Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc y): Performance tests are used to get results of HIP APIs.
There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn of process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf tests.
Sample Usage:
```cpp
hip::SpawnProc proc(<relative path of exe with test folder>, <optional bool value, if output is to be recorded>);
REQUIRE(0 == proc.run()); // Test of return value of the proc
REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process
```
The process can be a standalone exe (see tests/catch/unit/printfExe for more information).
General Guidelines:
# General Guidelines:
- Do not use the catch2 tags. Tags wont be used for filtering
- Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails)
- Check return of each HIP API and fail whenever there is a misma tch with hipSuccess or hiprtcSuccess.
tests/catch/external/Catch2/cmake/Catch2/Catch.cmake
Sağlanmış
+10
Dosyayı Görüntüle
@@ -203,10 +203,20 @@ function(hip_add_exe_to_target)
"${list_args}"
)
# Create shared lib of all tests
if(NOT RTC_TESTING)
add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object> $<TARGET_OBJECTS:KERNELS>)
else ()
add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $<TARGET_OBJECTS:Main_Object>)
if(HIP_PLATFORM STREQUAL "amd")
target_link_libraries(${_NAME} hiprtc)
else()
target_link_libraries(${_NAME} nvrtc)
endif()
endif()
catch_discover_tests(${_NAME} PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
if(UNIX)
set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs)
set(_LINKER_LIBS ${_LINKER_LIBS} -ldl)
else()
# res files are built resource files using rc files.
# use llvm-rc exe to build the res files
tests/catch/hipTestMain/config/config_amd_windows.json
+12 -34
Dosyayı Görüntüle
@@ -30,35 +30,9 @@
"Unit_hipStreamPerThread_DeviceReset_1",
"Unit_hipStreamPerThread_DeviceReset_2",
"Unit_hipManagedKeyword_MultiGpu",
"Unit_hipGraphAddDependencies_Functional",
"Unit_hipGraph_BasicFunctional",
"Unit_hipGraphClone_Functional",
"Unit_hipGraphClone_MultiThreaded",
"Unit_hipGraphAddHostNode_ClonedGraphwithHostNode",
"Unit_hipGraphAddHostNode_VectorSquare",
"Unit_hipGraphAddHostNode_BasicFunc",
"Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory",
"Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory",
"Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel",
"Unit_hipGraphExecHostNodeSetParams_Negative",
"Unit_hipGraphExecHostNodeSetParams_ClonedGraphwithHostNode",
"Unit_hipGraphExecHostNodeSetParams_BasicFunc",
"Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory",
"Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory",
"Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel",
"Unit_hipGraphDestroyNode_DestroyDependencyNode",
"Unit_hipGraphGetNodes_Functional",
"Unit_hipGraphGetRootNodes_Functional",
"Unit_hipGraphHostNodeSetParams_ClonedGraphwithHostNode",
"Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph",
"Unit_hipGraphAddChildGraphNode_SingleChildNode",
"Unit_hipGraphExecMemcpyNodeSetParams1D_Functional",
"Unit_hipGraphRemoveDependencies_ChangeComputeFunc",
"Unit_hipGraphExecUpdate_Negative_CountDiffer",
"Unit_hipGraphExecUpdate_Functional",
"Unit_hipGraphExecEventRecordNodeSetEvent_Negative",
"Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative",
"Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Functional",
"Unit_hipPtrGetAttribute_Simple",
"Unit_hipStreamCreateWithPriority_ValidateWithEvents",
"Unit_hipEvent",
@@ -67,20 +41,24 @@
"Unit_hipHostMalloc_Default",
"Unit_hipStreamCreate_MultistreamBasicFunctionalities",
"Unit_hipEventIpc",
"Unit_hipGraphAddDependencies_NegTest",
"Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
"Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime",
"Unit_hipGraphAddEventRecordNode_Functional_WithFlags",
"Unit_hipGraphAddEventRecordNode_MultipleRun",
"Unit_hipMalloc3D_Negative",
"Unit_hipPointerGetAttribute_MappedMem",
"Unit_hipGraphAddMemcpyNode1D_Functional",
"Unit_hipGraphAddMemcpyNode1D_Negative",
"Unit_hipStreamBeginCapture_BasicFunctional",
"Unit_hipStreamBeginCapture_hipStreamPerThread",
"Unit_hiprtc_functional",
"Unit_hipStreamValue_Write",
"Unit_hipStreamGetCaptureInfo_BasicFunctional",
"Unit_hipStreamGetCaptureInfo_hipStreamPerThread",
"Unit_hipStreamGetCaptureInfo_UniqueID",
"Unit_hipStreamGetCaptureInfo_ArgValidation",
"# Following test is related to ticket EXSWCPHIPT-41",
"Unit_hipStreamGetPriority_happy"
"Unit_hipStreamGetPriority_happy",
"Unit_hipMemPoolApi_Basic",
"Unit_hipMemPoolApi_BasicAlloc",
"Unit_hipMemPoolApi_BasicTrim",
"Unit_hipMemPoolApi_BasicReuse",
"Unit_hipMemPoolApi_Opportunistic",
"Unit_hipMemPoolApi_Default",
"Unit_hipDeviceGetUuid"
]
}
tests/catch/hipTestMain/hip_test_context.cc
+75 -16
Dosyayı Görüntüle
@@ -23,10 +23,10 @@ void TestContext::detectPlatform() {
std::string TestContext::substringFound(std::vector<std::string> list, std::string filename) {
std::string match = "";
for(unsigned int i = 0; i < list.size() ; i++) {
for (unsigned int i = 0; i < list.size(); i++) {
if (filename.find(list.at(i)) != std::string::npos) {
match = list.at(i);
break;
match = list.at(i);
break;
}
}
return match;
@@ -35,20 +35,19 @@ std::string TestContext::substringFound(std::vector<std::string> list, std::stri
std::string TestContext::getMatchingConfigFile(std::string config_dir) {
std::string configFileToUse;
for(auto& p: fs::recursive_directory_iterator(config_dir)) {
for (auto& p : fs::recursive_directory_iterator(config_dir)) {
fs::path filename = p.path();
std::string cur_arch = "TODO";
std::string arch = substringFound(amd_arch_list_,filename.filename().string());
std::string platform = substringFound(platform_list_,filename.filename().string());
std::string os = substringFound(os_list_,filename.filename().string());
std::string arch = substringFound(amd_arch_list_, filename.filename().string());
std::string platform = substringFound(platform_list_, filename.filename().string());
std::string os = substringFound(os_list_, filename.filename().string());
// if arch found then use that exit from loop
if(arch == cur_arch) {
if (arch == cur_arch) {
configFileToUse = filename.string();
break;
// match the platform/os and continue to look
} else if((platform == config_.platform) &&
(os == config_.os || os == "all")) {
configFileToUse = filename.string();
// match the platform/os and continue to look
} else if ((platform == config_.platform) && (os == config_.os || os == "all")) {
configFileToUse = filename.string();
}
}
return configFileToUse;
@@ -60,10 +59,10 @@ std::string& TestContext::getJsonFile() {
config_dir = config_dir.parent_path();
int levels = 0;
bool configFolderFound = false;
std::vector <std::string> configList;
std::vector<std::string> configList;
std::string configFile;
// check a max of 5 levels down the executable path
while(levels < 5) {
while (levels < 5) {
fs::path temp_path = config_dir;
temp_path /= "hipTestMain";
temp_path /= "config";
@@ -185,7 +184,7 @@ bool TestContext::parseJsonFile() {
return false;
}
const picojson::object &o = v.get<picojson::object>();
const picojson::object& o = v.get<picojson::object>();
for (picojson::object::const_iterator i = o.begin(); i != o.end(); ++i) {
// Processing for DisabledTests
if (i->first == "DisabledTests") {
@@ -196,7 +195,7 @@ bool TestContext::parseJsonFile() {
for (auto ai = val.begin(); ai != val.end(); ai++) {
std::string tmp = ai->get<std::string>();
std::string newRegexName;
for(const auto &c : tmp) {
for (const auto& c : tmp) {
if (c == '*')
newRegexName += ".*";
else
@@ -209,3 +208,63 @@ bool TestContext::parseJsonFile() {
return true;
}
void TestContext::cleanContext() {
for (auto& pair : compiledKernels) {
REQUIRE(hipSuccess == hipModuleUnload(pair.second.module));
}
}
void TestContext::trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule,
hipFunction_t kernelFunction) {
rtcState state{loadedModule, kernelFunction};
compiledKernels[kernelNameExpression] = state;
}
hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) {
auto it{compiledKernels.find(kernelNameExpression)};
if (it != compiledKernels.end()) {
return it->second.kernelFunction;
} else {
return nullptr;
}
}
void TestContext::addResults(HCResult r) {
std::unique_lock<std::mutex> lock(resultMutex);
results.push_back(r);
if ((!r.conditionsResult) ||
((r.result != hipSuccess) && (r.result != hipErrorPeerAccessAlreadyEnabled))) {
hasErrorOccured_.store(true);
}
}
void TestContext::finalizeResults() {
std::unique_lock<std::mutex> lock(resultMutex);
// clear the results whatever happens
std::shared_ptr<void> emptyVec(nullptr, [this](auto) { results.clear(); });
for (const auto& i : results) {
INFO("HIP API Result check\n File:: "
<< i.file << "\n Line:: " << i.line << "\n API:: " << i.call
<< "\n Result:: " << i.result << "\n Result Str:: " << hipGetErrorString(i.result));
REQUIRE(((i.result == hipSuccess) || (i.result == hipErrorPeerAccessAlreadyEnabled)));
REQUIRE(i.conditionsResult);
}
hasErrorOccured_.store(false); // Clear the flag
}
bool TestContext::hasErrorOccured() { return hasErrorOccured_.load(); }
TestContext::~TestContext() {
// Show this message when there are unchecked results
if (results.size() != 0) {
std::cerr << "HIP_CHECK_THREAD_FINALIZE() has not been called after HIP_CHECK_THREAD\n"
<< "Please call HIP_CHECK_THREAD_FINALIZE after joining threads\n"
<< "There is/are " << results.size() << " unchecked results from threads."
<< std::endl;
std::abort(); // Crash to bring users attention to this message and avoid accidental passing of
// tests without checking for errors
}
}
tests/catch/include/hip_test_checkers.hh
+141 -43
Dosyayı Görüntüle
@@ -23,17 +23,17 @@ THE SOFTWARE.
#pragma once
#include "hip_test_common.hh"
#include <iostream>
#include<fstream>
#include<regex>
#include <fstream>
#include <regex>
#include <type_traits>
#define guarantee(cond, str) \
{ \
if (!(cond)) { \
INFO("guarantee failed: " << str); \
abort(); \
} \
}
#define guarantee(cond, str) \
{ \
if (!(cond)) { \
INFO("guarantee failed: " << str); \
abort(); \
} \
}
namespace HipTest {
@@ -73,15 +73,15 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM
return mismatchCount;
}
template<typename T> // pointer type
bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) {
template <typename T> // pointer type
bool checkArray(T* hData, T* hOutputData, size_t width, size_t height, size_t depth = 1) {
for (size_t i = 0; i < depth; i++) {
for (size_t j = 0; j < height; j++) {
for (size_t k = 0; k < width; k++) {
int offset = i*width*height + j*width + k;
int offset = i * width * height + j * width + k;
if (hData[offset] != hOutputData[offset]) {
INFO("Mismatch at [" << i << "," << j << "," << k << "]:"
<< hData[offset] << "----" << hOutputData[offset]);
INFO("Mismatch at [" << i << "," << j << "," << k << "]:" << hData[offset] << "----"
<< hOutputData[offset]);
CHECK(false);
return false;
}
@@ -120,7 +120,7 @@ template <typename T> void setDefaultData(size_t numElements, T* A_h, T* B_h, T*
if (A_h) A_h[i] = 3;
if (B_h) B_h[i] = 4;
if (C_h) C_h[i] = 5;
} else if(std::is_same<T, char>::value || std::is_same<T, unsigned char>::value) {
} else if (std::is_same<T, char>::value || std::is_same<T, unsigned char>::value) {
if (A_h) A_h[i] = 'a';
if (B_h) B_h[i] = 'b';
if (C_h) C_h[i] = 'c';
@@ -185,6 +185,110 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost);
}
// Threaded version of setDefaultData to be called from multi thread tests
// Call HIP_CHECK_THREAD_FINALIZE after joining
template <typename T> void setDefaultDataT(size_t numElements, T* A_h, T* B_h, T* C_h) {
// Initialize the host data:
for (size_t i = 0; i < numElements; i++) {
if (std::is_same<T, int>::value || std::is_same<T, unsigned int>::value) {
if (A_h) A_h[i] = 3;
if (B_h) B_h[i] = 4;
if (C_h) C_h[i] = 5;
} else if (std::is_same<T, char>::value || std::is_same<T, unsigned char>::value) {
if (A_h) A_h[i] = 'a';
if (B_h) B_h[i] = 'b';
if (C_h) C_h[i] = 'c';
} else {
if (A_h) A_h[i] = 3.146f + i;
if (B_h) B_h[i] = 1.618f + i;
if (C_h) C_h[i] = 1.4f + i;
}
}
}
// Threaded version of initArraysForHost to be called from multi thread tests
// Call HIP_CHECK_THREAD_FINALIZE after joining
template <typename T>
void initArraysForHostT(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = false) {
size_t Nbytes = N * sizeof(T);
if (usePinnedHost) {
if (A_h) {
HIP_CHECK_THREAD(hipHostMalloc((void**)A_h, Nbytes));
}
if (B_h) {
HIP_CHECK_THREAD(hipHostMalloc((void**)B_h, Nbytes));
}
if (C_h) {
HIP_CHECK_THREAD(hipHostMalloc((void**)C_h, Nbytes));
}
} else {
if (A_h) {
*A_h = (T*)malloc(Nbytes);
REQUIRE_THREAD(*A_h != nullptr);
}
if (B_h) {
*B_h = (T*)malloc(Nbytes);
REQUIRE_THREAD(*B_h != nullptr);
}
if (C_h) {
*C_h = (T*)malloc(Nbytes);
REQUIRE_THREAD(*C_h != nullptr);
}
}
setDefaultDataT(N, A_h ? *A_h : nullptr, B_h ? *B_h : nullptr, C_h ? *C_h : nullptr);
}
// Threaded version of initArrays to be called from multi thread tests
// Call HIP_CHECK_THREAD_FINALIZE after joining
template <typename T>
void initArraysT(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
bool usePinnedHost = false) {
size_t Nbytes = N * sizeof(T);
if (A_d) {
HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes));
}
if (B_d) {
HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes));
}
if (C_d) {
HIP_CHECK_THREAD(hipMalloc(C_d, Nbytes));
}
initArraysForHostT(A_h, B_h, C_h, N, usePinnedHost);
}
// Threaded version of freeArraysForHost to be called from multi thread tests
// Call HIP_CHECK_THREAD_FINALIZE after joining
template <typename T> void freeArraysForHostT(T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (usePinnedHost) {
if (A_h) {
HIP_CHECK_THREAD(hipHostFree(A_h));
}
if (B_h) {
HIP_CHECK_THREAD(hipHostFree(B_h));
}
if (C_h) {
HIP_CHECK_THREAD(hipHostFree(C_h));
}
} else {
if (A_h) {
free(A_h);
}
if (B_h) {
free(B_h);
}
if (C_h) {
free(C_h);
}
}
}
template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (usePinnedHost) {
if (A_h) {
@@ -210,6 +314,21 @@ template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin
return true;
}
template <typename T>
void freeArraysT(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (A_d) {
HIP_CHECK_THREAD(hipFree(A_d));
}
if (B_d) {
HIP_CHECK_THREAD(hipFree(B_d));
}
if (C_d) {
HIP_CHECK_THREAD(hipFree(C_d));
}
freeArraysForHostT(A_h, B_h, C_h, usePinnedHost);
}
template <typename T>
bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (A_d) {
@@ -226,20 +345,6 @@ bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo
}
template <typename T>
unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock,
size_t N) {
int device;
HIP_CHECK(hipGetDevice(&device));
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
unsigned blocks = props.multiProcessorCount * blocksPerCU;
if (blocks * threadsPerBlock > N) {
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
}
return blocks;
}
template<typename T>
static bool assemblyFile_Verification(std::string assemfilename, std::string inst) {
std::string filePath = "./catch/unit/deviceLib/";
bool result = false;
@@ -254,34 +359,27 @@ static bool assemblyFile_Verification(std::string assemfilename, std::string ins
while (getline(file, line)) {
line_pos++;
if ((std::is_same<T, float>::value)) {
if (!start_pos &&
std::regex_search(line,
std::regex("Begin function (.*)AtomicCheck"))) {
if (!start_pos && std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) {
start_pos = line_pos;
}
if (!last_pos &&
std::regex_search(line,
std::regex(".Lfunc_end0-(.*)AtomicCheck"))) {
if (!last_pos && std::regex_search(line, std::regex(".Lfunc_end0-(.*)AtomicCheck"))) {
last_pos = line_pos;
break;
}
} else {
if ((start_match != 2) && std::regex_search(line,
std::regex("Begin function (.*)AtomicCheck"))) {
if ((start_match != 2) &&
std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) {
start_match++;
if (start_match == 2)
start_pos = line_pos;
if (start_match == 2) start_pos = line_pos;
}
if (!last_pos && std::regex_search(line,
std::regex("func_end1-(.*)AtomicCheck"))) {
if (!last_pos && std::regex_search(line, std::regex("func_end1-(.*)AtomicCheck"))) {
last_pos = line_pos;
break;
}
}
if (start_pos) {
result = std::regex_search(line, std::regex(inst));
if (result)
break;
if (result) break;
}
}
} else {
tests/catch/include/hip_test_common.hh
+142 -40
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021 - 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
@@ -23,7 +23,13 @@ THE SOFTWARE.
#pragma once
#include "hip_test_context.hh"
#include <catch.hpp>
#include <atomic>
#include <chrono>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <mutex>
#include <cstdlib>
#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
@@ -32,22 +38,51 @@ THE SOFTWARE.
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
INFO("Error: " << hipGetErrorString(localError) << " Code: " << localError << " Str: " \
<< #error << " In File: " << __FILE__ << " At line: " << __LINE__); \
INFO("Error: " << hipGetErrorString(localError) << "\n Code: " << localError \
<< "\n Str: " << #error << "\n In File: " << __FILE__ \
<< "\n At line: " << __LINE__); \
REQUIRE(false); \
} \
}
// Threaded HIP_CHECKs
#define HIP_CHECK_THREAD(error) \
{ \
/*To see if error has occured in previous threads, stop execution */ \
if (TestContext::get().hasErrorOccured() == true) { \
return; /*This will only work with std::thread and not with std::async*/ \
} \
auto localError = error; \
HCResult result(__LINE__, __FILE__, localError, #error); \
TestContext::get().addResults(result); \
}
#define REQUIRE_THREAD(condition) \
{ \
/*To see if error has occured in previous threads, stop execution */ \
if (TestContext::get().hasErrorOccured() == true) { \
return; /*This will only work with std::thread and not with std::async*/ \
} \
auto localResult = (condition); \
HCResult result(__LINE__, __FILE__, hipSuccess, #condition, localResult); \
TestContext::get().addResults(result); \
}
// Do not call before all threads have joined
#define HIP_CHECK_THREAD_FINALIZE() \
{ TestContext::get().finalizeResults(); }
// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError.
#define HIP_CHECK_ERROR(errorExpr, expectedError) \
{ \
hipError_t localError = errorExpr; \
INFO("Matching Errors: " \
<< " Expected Error: " << hipGetErrorString(expectedError) \
<< " Expected Code: " << expectedError << '\n' \
<< "\n Expected Error: " << hipGetErrorString(expectedError) \
<< "\n Expected Code: " << expectedError << '\n' \
<< " Actual Error: " << hipGetErrorString(localError) \
<< " Actual Code: " << localError << "\nStr: " << #errorExpr \
<< "\nIn File: " << __FILE__ << " At line: " << __LINE__); \
<< "\n Actual Code: " << localError << "\nStr: " << #errorExpr \
<< "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \
REQUIRE(localError == expectedError); \
}
@@ -56,8 +91,9 @@ THE SOFTWARE.
{ \
auto localError = error; \
if (localError != HIPRTC_SUCCESS) { \
INFO("Error: " << hiprtcGetErrorString(localError) << " Code: " << localError << " Str: " \
<< #error << " In File: " << __FILE__ << " At line: " << __LINE__); \
INFO("Error: " << hiprtcGetErrorString(localError) << "\n Code: " << localError \
<< "\n Str: " << #error << "\n In File: " << __FILE__ \
<< "\n At line: " << __LINE__); \
REQUIRE(false); \
} \
}
@@ -66,12 +102,6 @@ THE SOFTWARE.
#define HIP_ASSERT(x) \
{ REQUIRE((x)); }
#ifdef __cplusplus
#include <iostream>
#include <iomanip>
#include <chrono>
#endif
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
@@ -83,19 +113,20 @@ THE SOFTWARE.
}
#define HIPASSERT(condition) \
if (!(condition)) { \
printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \
abort(); \
}
if (!(condition)) { \
printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \
abort(); \
}
#if HT_NVIDIA
#define CTX_CREATE() \
hipCtx_t context;\
#define CTX_CREATE() \
hipCtx_t context; \
initHipCtx(&context);
#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context));
#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array));
#define HIP_TEX_REFERENCE hipTexRef
#define HIP_ARRAY hiparray
static void initHipCtx(hipCtx_t *pcontext) {
static void initHipCtx(hipCtx_t* pcontext) {
HIPCHECK(hipInit(0));
hipDevice_t device;
HIPCHECK(hipDeviceGet(&device, 0));
@@ -129,9 +160,9 @@ static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) {
}
static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) {
int device;
int device{0};
HIP_CHECK(hipGetDevice(&device));
hipDeviceProp_t props;
hipDeviceProp_t props{};
HIP_CHECK(hipGetDeviceProperties(&props, device));
unsigned blocks = props.multiProcessorCount * blocksPerCU;
@@ -142,23 +173,40 @@ static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlo
return blocks;
}
static inline int RAND_R(unsigned* rand_seed)
{
#if defined(_WIN32) || defined(_WIN64)
srand(*rand_seed);
return rand();
#else
return rand_r(rand_seed);
#endif
// Threaded version of setNumBlocks - to be used in multi threaded test
// Why? because catch2 does not support multithreaded macro calls
// Make sure you call HIP_CHECK_THREAD_FINALIZE after your threads join
// Also you can not return in threaded functions, due to how HIP_CHECK_THREAD works
static inline void setNumBlocksThread(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N,
unsigned& blocks) {
int device{0};
blocks = 0; // incase error has occured in some other thread and the next call might not execute,
// we set the blocks size to 0
HIP_CHECK_THREAD(hipGetDevice(&device));
hipDeviceProp_t props{};
HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device));
blocks = props.multiProcessorCount * blocksPerCU;
if (blocks * threadsPerBlock > N) {
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
}
}
static inline int RAND_R(unsigned* rand_seed) {
#if defined(_WIN32) || defined(_WIN64)
srand(*rand_seed);
return rand();
#else
return rand_r(rand_seed);
#endif
}
inline bool isImageSupported() {
int imageSupport = 1;
int imageSupport = 1;
#ifdef __HIP_PLATFORM_AMD__
int device;
HIP_CHECK(hipGetDevice(&device));
HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
device));
int device;
HIP_CHECK(hipGetDevice(&device));
HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device));
#endif
return imageSupport != 0;
}
@@ -171,11 +219,65 @@ static inline void HIP_SKIP_TEST(char const* const reason) noexcept {
// ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test.
std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl;
}
/**
* @brief Helper template that returns the expected arguments of a kernel.
*
* @return constexpr std::tuple<FArgs...> the expected arguments of the kernel.
*/
template <typename... FArgs> std::tuple<FArgs...> getExpectedArgs(void(FArgs...)){};
/**
* @brief Asserts that the types of the arguments of a function match exactly with the types in the
* function signature.
* This is necessary because HIP RTC does not do implicit casting of the kernel
* parameters.
* In order to get the kernel function signature, this function should only called when
* RTC is disabled.
*
* @tparam F the kernel function
* @tparam Args the parameters that will be passed to the kernel.
*/
template <typename F, typename... Args> void validateArguments(F f, Args...) {
using expectedArgsTuple = decltype(getExpectedArgs(f));
static_assert(std::is_same<expectedArgsTuple, std::tuple<Args...>>::value,
"Kernel arguments types must match exactly!");
}
/**
* @brief Launch a kernel using either HIP or HIP RTC.
*
* @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a
* template).
* @tparam K The kernel type. Expects a function or template when RTC is disabled. Expects a
* function pointer instead when RTC is enabled.
* @tparam Dim Can be either dim3 or int.
* @tparam Args A list of kernel arguments to be forwarded.
* @param kernel The kernel to be launched (defined in kernels.hh)
* @param numBlocks
* @param numThreads
* @param memPerBlock
* @param stream
* @param packedArgs A list of kernel arguments to be forwarded.
*/
template <typename... Typenames, typename K, typename Dim, typename... Args>
void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock,
hipStream_t stream, Args&&... packedArgs) {
#ifndef RTC_TESTING
validateArguments(kernel, packedArgs...);
kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
#else
launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
std::forward<Args>(packedArgs)...);
#endif
}
} // namespace HipTest
// This must be called in the beginning of image test app's main() to indicate whether image
// is supported.
#define checkImageSupport() \
if (!HipTest::isImageSupported()) \
{ printf("Texture is not support on the device. Skipped.\n"); return; }
#define checkImageSupport() \
if (!HipTest::isImageSupported()) { \
printf("Texture is not support on the device. Skipped.\n"); \
return; \
}
tests/catch/include/hip_test_context.hh
+75 -10
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -22,9 +22,15 @@ THE SOFTWARE.
#pragma once
#include <hip/hip_runtime.h>
#include <hip/hiprtc.h>
#include <atomic>
#include <mutex>
#include <vector>
#include <iostream>
#include <string>
#include <set>
#include <unordered_map>
#if defined(_WIN32)
#define HT_WIN 1
@@ -57,11 +63,23 @@ static int _log_enable = (std::getenv("HT_LOG_ENABLE") ? 1 : 0);
}
typedef struct Config_ {
std::string json_file; // Json file
std::string platform; // amd/nvidia
std::string os; // windows/linux
std::string json_file; // Json file
std::string platform; // amd/nvidia
std::string os; // windows/linux
} Config;
// Store Multi threaded results
struct HCResult {
size_t line; // Line of check (HIP_CHECK_THREAD or REQUIRE_THREAD)
std::string file; // File name of the check
hipError_t result; // hipResult for HIP_CHECK_THREAD, for conditions its hipSuccess
std::string call; // Call of HIP API or a bool condition
bool conditionsResult; // If bool condition, result of call. For HIP Calls its true
HCResult(size_t l, std::string f, hipError_t r, std::string c, bool b = true)
: line(l), file(f), result(r), call(c), conditionsResult(b) {}
};
class TestContext {
bool p_windows = false, p_linux = false; // OS
bool amd = false, nvidia = false; // HIP Platform
@@ -69,14 +87,20 @@ class TestContext {
std::string current_test;
std::set<std::string> skip_test;
std::string json_file_;
std::vector<std::string> platform_list_ = {"amd" , "nvidia"};
std::vector<std::string> os_list_ = {"windows", "linux", "all"};
std::vector<std::string> amd_arch_list_ = {};
std::vector<std::string> platform_list_ = {"amd", "nvidia"};
std::vector<std::string> os_list_ = {"windows", "linux", "all"};
std::vector<std::string> amd_arch_list_ = {};
struct rtcState {
hipModule_t module;
hipFunction_t kernelFunction;
};
std::unordered_map<std::string, rtcState> compiledKernels{};
Config config_;
std::string& getJsonFile();
std::string substringFound( std::vector<std::string> list,
std::string filename);
std::string substringFound(std::vector<std::string> list, std::string filename);
void detectOS();
void detectPlatform();
void fillConfig();
@@ -86,10 +110,16 @@ class TestContext {
std::string getMatchingConfigFile(std::string config_dir);
const Config& getConfig() const { return config_; }
TestContext(int argc, char** argv);
// Multi threaded checks helpers
std::mutex resultMutex;
std::vector<HCResult> results; // Multi threaded test results buffer
std::atomic<bool> hasErrorOccured_{false};
public:
static const TestContext& get(int argc = 0, char** argv = nullptr) {
static TestContext& get(int argc = 0, char** argv = nullptr) {
static TestContext instance(argc, argv);
return instance;
}
@@ -103,6 +133,41 @@ class TestContext {
const std::string& getCurrentTest() const { return current_test; }
std::string currentPath() const;
// Multi threaded results helpers
void addResults(HCResult r); // Add multi threaded results
void finalizeResults(); // Validate on all results
bool hasErrorOccured(); // Query if error has occured
/**
* @brief Unload all loaded modules.
* Note: This function needs to be called at the end of each test that uses RTC.
* It is not possible to unload the loaded modules without adding explicit code to the end
* of each test. This function exists only to provide a clean way to exit a test when using RTC.
* However, not unloading a module explicitly shouldn't have any effect on the outcome of
* the test.
*/
void cleanContext();
/**
* @brief Keeps track of all the already compiled rtc kernels.
*
* @param kernelNameExpression The name expression (e.g. hipTest::vectorADD<float>).
* @param loadedModule The loaded module.
* @param kernelFunction The hipFunction that will be used to run the kernel in the future.
*/
void trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule,
hipFunction_t kernelFunction);
/**
* @brief Get the already compiled hip rtc kernel function if it exists.
*
* @param kernelNameExpression The name expression (e.g. hipTest::vectorADD<float>).
* @return the hipFunction if it exists. nullptr otherwise
*/
hipFunction_t getFunction(const std::string kernelNameExpression);
TestContext(const TestContext&) = delete;
void operator=(const TestContext&) = delete;
~TestContext();
};
tests/catch/include/hip_test_helper.hh
+1 -2
Dosyayı Görüntüle
@@ -51,8 +51,7 @@ static size_t getMemoryAmount() {
#endif
}
static size_t getHostThreadCount(const size_t memPerThread,
const size_t maxThreads) {
static inline size_t getHostThreadCount(const size_t memPerThread, const size_t maxThreads) {
if (memPerThread == 0) return 0;
auto memAmount = getMemoryAmount();
const auto processor_count = std::thread::hardware_concurrency();
tests/catch/include/hip_test_kernels.hh
+8 -1
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -89,4 +89,11 @@ template <typename T> __global__ void vector_square(const T* A_d, T* C_d, size_t
}
}
template <typename T> __global__ void vector_cubic(const T* A_d, T* C_d, size_t N_ELMTS) {
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = gputhread; i < N_ELMTS; i += stride) {
C_d[i] = A_d[i] * A_d[i] * A_d[i];
}
}
} // namespace HipTest
tests/catch/include/hip_test_rtc.hh
+279
Dosyayı Görüntüle
@@ -0,0 +1,279 @@
/*
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.
*/
#pragma once
#include <hip/hip_runtime.h>
#include <hip/hiprtc.h>
#include <kernel_mapping.hh>
#include <catch.hpp>
#include <string>
#include <vector>
#include <iostream>
#include <fstream>
#include <sstream>
#include <set>
#include "hip/hip_runtime_api.h"
#include "hip_test_context.hh"
namespace HipTest {
struct KernelArgument {
const void* ptr;
size_t sizeRequirement;
size_t alignmentRequirement;
};
/**
* @brief Reconstructs the name expression for the kernel.
*
* @param kernelName the name of the kernel (e.g. "HipTest::VectorADD")
* @param typenames the typenames used by this kernel (e.g. "float").
* @return std::string the reconstructed expression (e.g. "VectorADD<float>""). Returns kernelName
* instead if the kernel is not a template.
*/
inline std::string reconstructExpression(std::string& kernelName,
std::vector<std::string>& typenames) {
std::string kernelExpression = kernelName;
if (typenames.size() > 0) {
kernelExpression += "<" + typenames[0];
for (size_t i = 1; i < typenames.size(); ++i) {
kernelExpression += "," + typenames[i];
}
kernelExpression += ">";
}
return kernelExpression;
}
/**
* @brief Packs the kernel arguments into the format expected by hipModuleLaunchKernel
*
* @param args list of arguments for the kernel and their alignemnt requirements.
* @return std::vector<char> the packed arguments ready to be passed on to hipModuleLaunchKernel
*/
inline std::vector<char> alignArguments(std::vector<KernelArgument>& args) {
std::vector<char> alignedArguments{};
int count = 0;
for (auto& arg : args) {
const char* argPtr{reinterpret_cast<const char*>(arg.ptr)};
/*
* Details about the padding formula can be found at:
* https://en.wikipedia.org/wiki/Data_structure_alignment#Data_structure_padding
*/
int paddingNeeded = -count & (arg.alignmentRequirement - 1);
alignedArguments.insert(std::end(alignedArguments), paddingNeeded, 0);
count += paddingNeeded;
alignedArguments.insert(std::end(alignedArguments), argPtr, argPtr + arg.sizeRequirement);
count += arg.sizeRequirement;
}
return alignedArguments;
}
inline std::vector<char> getKernelCode(hiprtcProgram& rtcProgram) {
size_t codeSize;
REQUIRE(HIPRTC_SUCCESS == hiprtcGetCodeSize(rtcProgram, &codeSize));
std::vector<char> code(codeSize);
REQUIRE(HIPRTC_SUCCESS == hiprtcGetCode(rtcProgram, code.data()));
return code;
}
/**
* @brief Compiles a kernel using HIP RTC
*
* @param rtcKernel the name of the kernel to compile.
* @param kernelNameExpression the name expression to be added to the RTC program (e.g.
* HipTest::VectorADD<float>)
* @return hiprtcProgram the compiled rtc program.
*/
inline hiprtcProgram compileRTC(std::string& rtcKernel, std::string& kernelNameExpression) {
std::string fileName = mapKernelToFileName.at(rtcKernel);
std::string filePath{KERNELS_PATH + fileName};
INFO("Opening Kernel File: " << filePath);
std::ifstream kernelFile{filePath};
REQUIRE(kernelFile.is_open());
std::stringstream stringStream;
std::string line;
while (getline(kernelFile, line)) {
/* Skip the include directive since it is not part of the kernel */
if (line.find("#include") != std::string::npos) {
continue;
}
stringStream << line << '\n';
}
kernelFile.close();
std::string kernelCode{stringStream.str()};
INFO("RTC Kernel Code:\n" << kernelCode)
hiprtcProgram rtcProgram;
hiprtcCreateProgram(&rtcProgram, kernelCode.c_str(), (fileName + ".cu").c_str(), 0, nullptr,
nullptr);
std::vector<const char*> options{};
#ifdef __HIP_PLATFORM_AMD__
int deviceCount;
REQUIRE(hipSuccess == hipGetDeviceCount(&deviceCount));
std::set<std::string> architectures{};
for (int i = 0; i < deviceCount; ++i) {
hipDeviceProp_t props;
REQUIRE(hipSuccess == hipGetDeviceProperties(&props, i));
architectures.insert(std::string{"--gpu-architecture="} + props.gcnArchName);
}
for (auto& architecture : architectures) {
options.push_back(architecture.c_str());
}
#else
options.push_back("--fmad=false");
#endif
REQUIRE(HIPRTC_SUCCESS == hiprtcAddNameExpression(rtcProgram, kernelNameExpression.c_str()));
REQUIRE(HIPRTC_SUCCESS == hiprtcCompileProgram(rtcProgram, 1, options.data()));
return rtcProgram;
}
/**
* @brief Get a typename as a string
*
* @tparam T The typename
* @return std::string the string representation of T
*/
template <typename T> std::string getTypeName() {
std::string name, prefix, suffix;
#ifdef __clang__
name = __PRETTY_FUNCTION__;
prefix = "std::string HipTest::getTypeName() [T = ";
suffix = "]";
#elif defined(__GNUC__)
name = __PRETTY_FUNCTION__;
prefix = "std::string HipTest::getTypeName() [with T = ";
suffix = "; std::string = std::__cxx11::basic_string<char>]";
#elif defined(_MSC_VER)
name = __FUNCSIG__;
prefix = "std::string __cdecl HipTest::getTypeName<";
suffix = ">(void)";
#endif
return name.substr(prefix.size(), name.rfind(suffix) - prefix.size());
}
/**
* @brief Tells the user that the kernels are using HIP RTC. Prints only once per test.
*
*/
static inline void printInfo() {
static bool alreadyPrinted{false};
if (!alreadyPrinted) {
std::cout << "INFO: This test is running using HIP RTC to compile and run the kernels."
<< std::endl;
alreadyPrinted = true;
}
}
/**
* @brief Compiles and launches a kernel using HIP RTC
*
* @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a
* template).
* @tparam Args A list of kernel arguments to be forwarded.
* @param getKernelName A function wrapper that returns the name of the kernel to launch (check
* kernels.hh for more info)
* @param numBlocks
* @param numThreads
* @param memPerBlock
* @param stream
* @param packedArgs A list of kernel arguments to be forwarded.
*/
template <typename... Typenames, typename... Args>
void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThreads,
std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) {
printInfo();
TestContext& testContext = TestContext::get();
std::string kernelName = (*getKernelName)();
std::vector<std::string> kernelTypenames{std::string(HipTest::getTypeName<Typenames>())...};
std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames);
if (testContext.getFunction(kernelExpression) == nullptr) {
hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)};
std::vector<char> compiledCode{getKernelCode(rtcProgram)};
hipModule_t module;
REQUIRE(hipSuccess == hipModuleLoadData(&module, compiledCode.data()));
hipFunction_t kernelFunction;
const char* loweredName;
REQUIRE(HIPRTC_SUCCESS ==
hiprtcGetLoweredName(rtcProgram, kernelExpression.c_str(), &loweredName));
REQUIRE(hipSuccess == hipModuleGetFunction(&kernelFunction, module, loweredName));
/* After obtaining the kernelFunction, the program is no longer needed. So it can be destroyed */
REQUIRE(HIPRTC_SUCCESS == hiprtcDestroyProgram(&rtcProgram));
testContext.trackRtcState(kernelExpression, module, kernelFunction);
}
hipFunction_t kernelFunction = testContext.getFunction(kernelExpression);
std::vector<KernelArgument> args = {
{reinterpret_cast<const void*>(&packedArgs), sizeof(Args), alignof(Args)}...};
std::vector<char> alignedArguments{alignArguments(args)};
size_t argumentsSize{alignedArguments.size()};
void* config_array[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, alignedArguments.data(),
HIP_LAUNCH_PARAM_BUFFER_SIZE, reinterpret_cast<void*>(&argumentsSize),
HIP_LAUNCH_PARAM_END};
REQUIRE(hipSuccess ==
hipModuleLaunchKernel(kernelFunction, numBlocks.x, numBlocks.y, numBlocks.z, numThreads.x,
numThreads.y, numThreads.z, memPerBlock, stream, nullptr,
config_array));
}
/**
* @brief Template overload for when numBlocks and numThreads is an integer.
*
*/
template <typename... Typenames, typename... Args>
void launchRTCKernel(std::string kernelName, int numBlocks, int numThreads,
std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) {
launchRTCKernel<Typenames...>(kernelName, dim3(numBlocks), dim3(numThreads), memPerBlock, stream,
std::forward<Args>(packedArgs)...);
}
} // namespace HipTest
tests/catch/unit/memory/hipMemcpyToSymbol.cc → tests/catch/include/kernel_mapping.hh
+7 -23
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -17,27 +17,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define SIZE 1024
#pragma once
/* Test verifies hipMemcpyToSymbol API Negative scenarios.
*/
#include <map>
TEST_CASE("Unit_hipMemcpyToSymbol_Negative") {
void *Sd;
char S[SIZE]="This is not a device symbol";
HIP_CHECK(hipMalloc(&Sd, SIZE));
SECTION("Passing void pointer") {
REQUIRE(hipSuccess != hipMemcpyToSymbol(HIP_SYMBOL(Sd), S,
SIZE, 0, hipMemcpyDeviceToHost));
}
SECTION("Passing NULL Pointer") {
REQUIRE(hipSuccess != hipMemcpyToSymbol(nullptr, S,
SIZE, 0, hipMemcpyDeviceToHost));
}
HIP_CHECK(hipFree(Sd));
}
const std::map<std::string, std::string> mapKernelToFileName{
{"Set", "Set.cpp"},
{"HipTest::vectorADD", "vectorADD.inl"},
};
tests/catch/include/kernels.hh
+55
Dosyayı Görüntüle
@@ -0,0 +1,55 @@
/*
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.
*/
#pragma once
#include <hip_test_common.hh>
#include <map>
#ifndef RTC_TESTING
__global__ void Set(int* Ad, int val);
/* Kernel Templates */
#include "vectorADD.inl"
#else
/*
* Wrapper Macros that create a string representation of the kernel name.
* In the case of kernel templates, a variadic template is used to ensure compatibility with
* the launchKernel template when RTC is not enabled. If the kernel is inside a namespace, use the
* "_NS" version of the Macro.
*/
#define FUNCTION_WRAPPER(param) \
std::string param() { return #param; }
#define TEMPLATE_WRAPPER(param) \
template <typename...> std::string param() { return #param; }
#define FUNCTION_WRAPPER_NS(param, namespace) \
std::string param() { return #namespace "::" #param; }
#define TEMPLATE_WRAPPER_NS(param, namespace) \
template <typename...> std::string param() { return #namespace "::" #param; }
FUNCTION_WRAPPER(Set);
namespace HipTest {
TEMPLATE_WRAPPER_NS(vectorADD, HipTest);
}
#endif
tests/catch/kernels/CMakeLists.txt
+8
Dosyayı Görüntüle
@@ -0,0 +1,8 @@
if(NOT RTC_TESTING)
set(TEST_SRC
Set.cpp
)
add_library(KERNELS EXCLUDE_FROM_ALL OBJECT ${TEST_SRC})
target_compile_options(KERNELS PUBLIC -std=c++17)
endif()
tests/catch/kernels/Set.cpp
+6
Dosyayı Görüntüle
@@ -0,0 +1,6 @@
#include <kernels.hh>
__global__ void Set(int* Ad, int val) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tx] = val;
}
tests/catch/kernels/vectorADD.inl
+10
Dosyayı Görüntüle
@@ -0,0 +1,10 @@
namespace HipTest {
template <typename T> __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < NELEM; i += stride) {
C_d[i] = A_d[i] + B_d[i];
}
}
}
tests/catch/stress/memory/CMakeLists.txt
+1
Dosyayı Görüntüle
@@ -4,6 +4,7 @@ set(TEST_SRC
hipMemcpyMThreadMSize.cc
hipMallocManagedStress.cc
hipMemPrftchAsyncStressTst.cc
hipHostMalloc.cc
)
hip_add_exe_to_target(NAME memory
tests/catch/stress/memory/hipHostMalloc.cc
+52
Dosyayı Görüntüle
@@ -0,0 +1,52 @@
/*
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.
*/
#include "hip_test_common.hh"
#include "hip_test_helper.hh"
// Stress allocation tests
// Try to allocate as much memory as possible
// But since max allocation can fail, we need to try the next value
TEST_CASE("Stress_hipHostMalloc_MaxAllocation") {
size_t devMemAvail{0}, devMemFree{0};
HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail));
auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes
REQUIRE(devMemFree > 0);
REQUIRE(devMemAvail > 0);
REQUIRE(hostMemFree > 0);
size_t memFree = std::min(devMemFree, hostMemFree); // which is the limiter cpu or gpu
char* d_ptr{nullptr};
size_t counter{0};
INFO("Max Allocation of " << memFree << " bytes!");
while (hipHostMalloc(&d_ptr, memFree) != hipSuccess && memFree > 1) {
counter++;
INFO("Attempt to allocate " << memFree << " bytes out of " << devMemFree << "bytes Failed!");
memFree >>= 1; // reduce the memory to be allocated by half
REQUIRE(counter <= 2); // Make sure that we are atleast able to allocate 1/4th of max memory
}
HIP_CHECK(hipMemset(d_ptr, 1, memFree));
HIP_CHECK(hipDeviceSynchronize()); // Flush caches
REQUIRE(std::all_of(d_ptr, d_ptr + memFree, [](unsigned char n) { return n == 1; }));
HIP_CHECK(hipHostFree(d_ptr));
}
tests/catch/stress/memory/hipMallocManagedStress.cc
-1
Dosyayı Görüntüle
@@ -67,7 +67,6 @@ __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) {
Hmm[i] = Dptr[i] * 10;
}
}
static bool IfTestPassed = true;
static void LaunchKrnl4(size_t NumElms, int InitVal) {
int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0;
tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc
+2 -2
Dosyayı Görüntüle
@@ -26,8 +26,8 @@ THE SOFTWARE.
// Kernel function
__global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) {
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < N; i += stride) {
Hmm[i] = Hmm[i] * Hmm[i];
}
tests/catch/stress/printf/Stress_printf_ComplexKernels.cc
+3 -3
Dosyayı Görüntüle
@@ -230,19 +230,19 @@ __device__ __host__ struct printInfo startPrint(uint32_t tid,
// This kernel is launched only in X dimension
__global__ void kernel_complex_opX(uint32_t *a, uint32_t *b,
uint32_t iterCount) {
uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
startPrint(tid, iterCount, a, b);
}
// This kernel is launched only in Y dimension
__global__ void kernel_complex_opY(uint32_t *a, uint32_t *b,
uint32_t iterCount) {
uint32_t tid = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
uint32_t tid = threadIdx.y + blockIdx.y * blockDim.y;
startPrint(tid, iterCount, a, b);
}
// This kernel is launched only in Z dimension
__global__ void kernel_complex_opZ(uint32_t *a, uint32_t *b,
uint32_t iterCount) {
uint32_t tid = hipThreadIdx_z + hipBlockIdx_z * hipBlockDim_z;
uint32_t tid = threadIdx.z + blockIdx.z * blockDim.z;
startPrint(tid, iterCount, a, b);
}
#ifdef __linux__
tests/catch/stress/printf/Stress_printf_SimpleKernels.cc
+6 -6
Dosyayı Görüntüle
@@ -85,7 +85,7 @@ __global__ void kernel_printf_conststr(uint iterCount) {
// 'g' grid size such that (total bytes per iteration)*n*b*g ≈ N GB,
// where N is user input.
__global__ void kernel_printf_two_conditionalstr(uint iterCount) {
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint tid = threadIdx.x + blockIdx.x * blockDim.x;
uint mod_tid = (tid % 2);
if (0 == mod_tid) {
for (uint count = 0; count < iterCount; count++) {
@@ -101,7 +101,7 @@ __global__ void kernel_printf_two_conditionalstr(uint iterCount) {
// iterations per thread using 'b' block size and 'g' grid size such that
// (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
__global__ void kernel_printf_single_conditionalstr(uint iterCount) {
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint tid = threadIdx.x + blockIdx.x * blockDim.x;
uint mod_tid = (tid % 2);
if (0 == mod_tid) {
for (uint count = 0; count < iterCount; count++) {
@@ -115,7 +115,7 @@ __global__ void kernel_printf_single_conditionalstr(uint iterCount) {
// iterations per thread using 'b' block size and 'g' grid size such
// that (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
__global__ void kernel_printf_variablestr(uint iterCount, int *ret) {
uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint tid = threadIdx.x + blockIdx.x * blockDim.x;
int retlocal = 0;
const char *const_str =
"Hello World from Device.Iam printing (threadID,number)=";
@@ -134,7 +134,7 @@ __global__ void kernel_printf_variablestr(uint iterCount, int *ret) {
// size and 'g' grid size such that
// (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
__global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) {
uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
int retlocal = 0;
const char *const_str =
"Hello World from Device.Iam printing number=";
@@ -158,7 +158,7 @@ __global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) {
// (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
__global__ void kernel_dependent_calc_atomic(uint32_t iterCount,
int *ret) {
uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
int retlocal = 0;
const char *const_str =
"Hello World from Device.Iam printing number=";
@@ -209,7 +209,7 @@ __global__ void kernel_shared_mem() {
__shared__ uint32_t sharedMem;
sharedMem = 0;
__syncthreads();
atomicAdd(&sharedMem, hipThreadIdx_x);
atomicAdd(&sharedMem, threadIdx.x);
__syncthreads();
printf("%s%u\n", CONST_STR3, sharedMem);
}
tests/catch/stress/stream/CMakeLists.txt
+3 -1
Dosyayı Görüntüle
@@ -1,8 +1,10 @@
# Common Tests - Test independent of all platforms
set(TEST_SRC
Stress_hipStreamCreate.cc
streamEnqueue.cc
)
hip_add_exe_to_target(NAME stream
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME stress_test)
TEST_TARGET_NAME stress_test
COMPILE_OPTIONS -std=c++14)
tests/catch/stress/stream/streamEnqueue.cc
+231
Dosyayı Görüntüle
@@ -0,0 +1,231 @@
/*
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 <algorithm>
#include <atomic>
#include <map>
#include <mutex>
#include <random>
#include <thread>
__global__ void addVal(unsigned long long* ptr, size_t index, unsigned long long val) {
atomicAdd(ptr + index, val);
}
// Create a copy constructible AtomicWrap around std::atomic so that we can put it in a vector
template <typename T> struct AtomicWrap {
std::atomic<T> data;
AtomicWrap() : data() {}
AtomicWrap(T i) : data(i) {}
AtomicWrap(const std::atomic<T>& a) : data(a.load()) {}
AtomicWrap(const AtomicWrap& other) : data(other.data.load()) {}
AtomicWrap& operator=(const AtomicWrap& other) {
data.store(other.data.load());
return *this;
}
};
// Have multiple threads and enqueue commands from them on a single stream
// Validate at the end that all commands have completed successfully
TEST_CASE("Stress_StreamEnqueue_DifferentThreads") {
auto hwThreads = std::thread::hardware_concurrency();
hwThreads = (hwThreads >= 2) ? hwThreads : 2; // Run atleast 2 threads
std::vector<AtomicWrap<unsigned long long>> hostData(hwThreads, 0);
unsigned long long* dPtr{nullptr};
HIP_CHECK(hipMalloc(&dPtr, sizeof(unsigned long long) * hwThreads));
REQUIRE(dPtr != nullptr);
HIP_CHECK(hipMemset(dPtr, 0, sizeof(unsigned long long) * hwThreads));
std::random_device device;
std::mt19937 engine(device());
constexpr size_t maxWork = 10000;
constexpr size_t maxVal = 10;
std::uniform_int_distribution<std::mt19937::result_type> genIndex(0, hwThreads - 1);
std::uniform_int_distribution<std::mt19937::result_type> genWork(0, maxWork);
std::uniform_int_distribution<std::mt19937::result_type> genVal(0, maxVal);
auto enqueueKernelThread = [&](hipStream_t stream) {
auto iter = genWork(engine); // Generate work to be done via thread
for (auto i = 0; i < iter; i++) {
auto index = genIndex(engine); // Generate Index to add to
auto val = genVal(engine); // Generate value to add to the destination
hostData[index].data += val; // Replicate it on host
addVal<<<1, 1, 0, stream>>>(dPtr, static_cast<size_t>(index),
static_cast<unsigned long long>(val)); // And on device
}
};
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
std::vector<std::thread> threadPool{};
threadPool.reserve(hwThreads);
// Launch work
for (size_t i = 0; i < hwThreads; i++) {
threadPool.emplace_back(std::thread(enqueueKernelThread, stream));
}
// Wait for work to finish
for (auto& i : threadPool) {
i.join();
}
HIP_CHECK(hipStreamDestroy(stream));
auto hPtr = std::make_unique<unsigned long long[]>(hwThreads);
HIP_CHECK(
hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads, hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(dPtr));
// Validate that CPU and GPU has the same results
for (size_t i = 0; i < hwThreads; i++) {
INFO("Check for Index " << i);
REQUIRE(hostData[i].data.load() == hPtr[i]);
}
}
__global__ void doOperation(int* dPtr, size_t size, int val) {
auto i = threadIdx.x;
atomicAdd(dPtr + i, val);
}
// Allocate mulitple stream for same device.
// Same device stream operate on same memory
TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") {
int deviceCount{0};
HIP_CHECK(hipGetDeviceCount(&deviceCount));
REQUIRE(deviceCount > 0);
// Skip the test if devices less than 2
if (deviceCount <= 1) {
HipTest::HIP_SKIP_TEST("Skipping because devices <= 1");
return;
}
constexpr size_t streamPerGPU{3}; // Stream per gpu
std::vector<hipStream_t> streamPool{};
streamPool.reserve(deviceCount * streamPerGPU);
std::map<hipStream_t, int*> streamToDeviceMemory; // Map of stream and device memory
std::map<hipStream_t, AtomicWrap<int>> streamToHostMemory; // Map of stream and host result
std::map<hipStream_t, size_t> streamToDeviceIndex; // Map of stream and device it was created on
constexpr size_t size = 1024;
for (size_t i = 0; i < deviceCount; i++) {
HIP_CHECK(hipSetDevice(i));
for (size_t j = 0; j < streamPerGPU; j++) {
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
REQUIRE(stream != nullptr);
streamPool.push_back(stream);
int* dPtr{nullptr};
HIP_CHECK(hipMalloc(&dPtr, sizeof(int) * size));
REQUIRE(dPtr != nullptr);
HIP_CHECK(hipMemset(dPtr, 0, sizeof(int) * size));
streamToDeviceMemory[stream] = dPtr; // All streams work on exclusive memory
streamToHostMemory[stream] = AtomicWrap<int>(0); // CPU result
streamToDeviceIndex[stream] = i; // Capture device id for stream
}
}
constexpr size_t maxVal = 5;
constexpr size_t maxWorkPerThread = 10000;
// Boiler plate code to generate a random number
std::random_device device;
std::mt19937 engine(device());
std::uniform_int_distribution<std::mt19937::result_type> genVal(-maxVal, maxVal);
std::uniform_int_distribution<std::mt19937::result_type> genStream(0, streamPool.size() - 1);
#if HT_NVIDIA
std::mutex ness; // On nvidia, current device needs to match stream's device
#endif
auto enqueueKernelThread = [&]() {
for (size_t i = 0; i < maxWorkPerThread; i++) {
#if HT_NVIDIA
std::unique_lock<std::mutex> lock(ness); // Lock on creation
#endif
hipStream_t stream = streamPool[genStream(engine)]; // Get a random stream
// TODO use HIP_CHECK_THREAD when PR#2664 is merged
if (hipSuccess != hipSetDevice(streamToDeviceIndex[stream])) {
return;
}
int val = genVal(engine); // Generate Value to add/sub to
streamToHostMemory[stream].data.fetch_add(val); // Replicate result on CPU
auto dPtr = streamToDeviceMemory[stream];
doOperation<<<1, 1024, 0, stream>>>(dPtr, size,
val); // On GPU
}
};
auto maxThreads = std::thread::hardware_concurrency();
maxThreads = (maxThreads >= 2) ? maxThreads : 2; // Run atleast 2 threads
std::vector<std::thread> threadPool{};
threadPool.reserve(maxThreads);
// Launch Threads
for (size_t i = 0; i < maxThreads; i++) {
threadPool.emplace_back(std::thread(enqueueKernelThread));
}
// Wait for them to stop
for (auto& i : threadPool) {
i.join();
}
// Sync and check results
for (auto& i : streamPool) {
HIP_CHECK(hipStreamSynchronize(i));
auto dResult = std::make_unique<int[]>(size);
HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i], sizeof(int) * size,
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(streamToDeviceMemory[i]));
HIP_CHECK(hipStreamDestroy(i));
auto res = streamToHostMemory[i].data.load();
INFO("Matching CPU: " << res << " GPU: " << dResult[0] << " Dev Ptr: "
<< streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]);
REQUIRE(std::all_of(dResult.get(), dResult.get() + size, [=](int r) { return r == res; }));
}
}
tests/catch/unit/device/CMakeLists.txt
+1
Dosyayı Görüntüle
@@ -15,6 +15,7 @@ set(TEST_SRC
hipRuntimeGetVersion.cc
hipSetDeviceFlags.cc
hipSetGetDevice.cc
hipDeviceGetUuid.cc
)
hip_add_exe_to_target(NAME DeviceTest
tests/catch/unit/device/hipDeviceGetName.cc
+85 -17
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -21,49 +21,117 @@ THE SOFTWARE.
* Conformance test for checking functionality of
* hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
*/
#include <cstddef>
#include <hip_test_common.hh>
#include <cstring>
#include <cstdio>
#include <array>
#include <algorithm>
#include <iterator>
constexpr size_t LEN = 256;
#define LEN 256
/**
* hipDeviceGetName tests
* Scenario1: Validates the name string with hipDeviceProp_t.name[256]
* Scenario2: Validates returned error code for name = nullptr
* Scenario3: Validates returned error code for len = 0
* Scenario4: Validates returned error code for len < 0
* Scenario5: Validates returned error code for an invalid device
* Scenario6: Validates partially filling the name into a char array
*/
TEST_CASE("Unit_hipDeviceGetName-NegTst") {
TEST_CASE("Unit_hipDeviceGetName_NegTst") {
std::array<char, LEN> name;
int numDevices = 0;
char name[LEN];
hipDevice_t device;
HIP_CHECK(hipGetDeviceCount(&numDevices));
std::vector<hipDevice_t> devices(numDevices);
for (int i = 0; i < numDevices; i++) {
HIP_CHECK(hipDeviceGet(&device, i));
HIP_CHECK(hipDeviceGetName(name, LEN, device));
// Scenario2
CHECK_FALSE(hipSuccess == hipDeviceGetName(nullptr, LEN, device));
HIP_CHECK(hipDeviceGet(&devices[i], i));
}
SECTION("Valid Device") {
const auto device = GENERATE_COPY(from_range(std::begin(devices), std::end(devices)));
SECTION("Nullptr for name argument") {
// Scenario2
HIP_CHECK_ERROR(hipDeviceGetName(nullptr, name.size(), device), hipErrorInvalidValue);
}
#if HT_AMD
// These test scenarios fail on NVIDIA.
// Scenario3
CHECK_FALSE(hipSuccess == hipDeviceGetName(name, 0, device));
// Scenario4
CHECK_FALSE(hipSuccess == hipDeviceGetName(name, -1, device));
SECTION("Zero name length") {
// Scenario3
HIP_CHECK_ERROR(hipDeviceGetName(name.data(), 0, device), hipErrorInvalidValue);
}
SECTION("Negative name length") {
// Scenario4
HIP_CHECK_ERROR(hipDeviceGetName(name.data(), -1, device), hipErrorInvalidValue);
}
#endif
}
SECTION("Invalid Device") {
hipDevice_t badDevice = devices.back() + 1;
constexpr size_t timeout = 100;
size_t timeoutCount = 0;
while (std::find(std::begin(devices), std::end(devices), badDevice) != std::end(devices)) {
badDevice += 1;
timeoutCount += 1;
REQUIRE(timeoutCount < timeout); // give up after a while
}
// Scenario5
HIP_CHECK_ERROR(hipDeviceGetName(name.data(), name.size(), badDevice), hipErrorInvalidDevice);
}
}
TEST_CASE("Unit_hipDeviceGetName-CheckPropName") {
TEST_CASE("Unit_hipDeviceGetName_CheckPropName") {
int numDevices = 0;
char name[LEN];
std::array<char, LEN> name;
hipDevice_t device;
hipDeviceProp_t prop;
HIP_CHECK(hipGetDeviceCount(&numDevices));
for (int i = 0; i < numDevices; i++) {
HIP_CHECK(hipDeviceGet(&device, i));
HIP_CHECK(hipDeviceGetName(name, LEN, device));
HIP_CHECK(hipDeviceGetName(name.data(), name.size(), device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
// Scenario1
CHECK_FALSE(0 != strcmp(name, prop.name));
CHECK(strncmp(name.data(), prop.name, name.size()) == 0);
}
}
TEST_CASE("Unit_hipDeviceGetName_PartialFill") {
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-108");
return;
#endif
std::array<char, LEN> name;
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
auto ordinal = GENERATE_COPY(range(0, numDevices));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, ordinal));
HIP_CHECK(hipDeviceGetName(name.data(), name.size(), device));
auto start = std::begin(name);
auto end = std::end(name);
const auto len = std::distance(start, std::find(start, end, 0));
// fill up only half of the length
const auto fillLen = len / 2;
constexpr char fillValue = 1;
std::fill(start, end, fillValue);
// Scenario6
HIP_CHECK(hipDeviceGetName(name.data(), fillLen, device));
const auto strEnd = start + fillLen - 1;
REQUIRE(std::all_of(start, strEnd, [](char& c) { return c != 0; }));
REQUIRE(*strEnd == 0);
REQUIRE(std::all_of(strEnd+1, end, [](char& c) { return c == fillValue; }));
}
tests/catch/unit/device/hipDeviceGetUuid.cc
+51
Dosyayı Görüntüle
@@ -0,0 +1,51 @@
/*
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, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
* Conformance test for checking functionality of
* hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device);
*/
#include <hip_test_common.hh>
#include <cstring>
#include <cstdio>
/**
* hipDeviceGetUuid tests
* Scenario1: Validates the returned UUID
* Scenario2: Validates returned error code for UUID = nullptr
* Scenario3 & 4: Validates returned error code for invalid device
*/
TEST_CASE("Unit_hipDeviceGetUuid") {
int numDevices = 0;
hipDevice_t device;
hipUUID uuid;
HIP_CHECK(hipGetDeviceCount(&numDevices));
for (int i = 0; i < numDevices; i++) {
HIP_CHECK(hipDeviceGet(&device, i));
// Scenario 1
HIP_CHECK(hipDeviceGetUuid(&uuid, device));
REQUIRE_FALSE(!strcmp(uuid.bytes, ""));
// Scenario 2
REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(nullptr, device));
}
// Scenario 3
REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(&uuid, -1));
// Scenario 4
REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(&uuid, numDevices));
}
tests/catch/unit/event/Unit_hipEventRecord.cc
+5 -5
Dosyayı Görüntüle
@@ -23,8 +23,8 @@ THE SOFTWARE.
// Through manual inspection of the reported timestamps, can determine if recording a NULL event
// forces synchronization : set
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include <kernels.hh>
#include <hip_test_context.hh>
#include <hip_test_common.hh>
TEST_CASE("Unit_hipEventRecord") {
@@ -61,8 +61,8 @@ TEST_CASE("Unit_hipEventRecord") {
// Record the start event
HIP_CHECK(hipEventRecord(start, NULL));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
HipTest::launchKernel<float>(HipTest::vectorADD<float>, blocks, threadsPerBlock, 0, 0,
static_cast<const float*>(A_d), static_cast<const float*>(B_d), C_d, N);
HIP_CHECK(hipEventRecord(stop, NULL));
HIP_CHECK(hipEventSynchronize(stop));
@@ -87,5 +87,5 @@ TEST_CASE("Unit_hipEventRecord") {
HIP_CHECK(hipEventDestroy(stop));
HipTest::checkVectorADD(A_h, B_h, C_h, N, true);
TestContext::get().cleanContext();
}
tests/catch/unit/event/Unit_hipEvent_Negative.cc
+32 -1
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -46,3 +46,34 @@ TEST_CASE("Unit_hipEventDestroy_NullCheck") {
auto res = hipEventDestroy(nullptr);
REQUIRE(res != hipSuccess);
}
TEST_CASE("Unit_hipEventCreate_IncompatibleFlags") {
hipEvent_t event;
#if HT_AMD
HipTest::HIP_SKIP_TEST("EXSWCPHIPT-106");
return;
#endif
HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, hipEventInterprocess), hipErrorInvalidValue);
#if HT_AMD
HIP_CHECK_ERROR(
hipEventCreateWithFlags(&event, hipEventReleaseToDevice | hipEventReleaseToSystem),
hipErrorInvalidValue);
#endif
unsigned allFlags{hipEventReleaseToDevice | hipEventReleaseToSystem | hipEventBlockingSync |
hipEventDisableTiming | hipEventDefault | hipEventInterprocess};
#if HT_AMD
HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, allFlags), hipErrorInvalidValue);
#else
/* Works on Non-AMD because hipEventReleaseToDevice / hipEventReleaseToSystem have no meaning in
* that case */
HIP_CHECK(hipEventCreateWithFlags(&event, allFlags));
#endif
unsigned invalidFlag{0x08000000};
HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, invalidFlag), hipErrorInvalidValue);
}
tests/catch/unit/graph/CMakeLists.txt
+4
Dosyayı Görüntüle
@@ -49,6 +49,7 @@ set(TEST_SRC
hipGraphInstantiate.cc
hipGraphExecUpdate.cc
hipGraphExecEventRecordNodeSetEvent.cc
hipGraphEventWaitNodeSetEvent.cc
hipGraphMemsetNodeGetParams.cc
hipGraphMemsetNodeSetParams.cc
hipGraphExecMemcpyNodeSetParamsFromSymbol.cc
@@ -58,6 +59,9 @@ set(TEST_SRC
hipGraphExecMemcpyNodeSetParams.cc
hipStreamBeginCapture.cc
hipGraphAddMemcpyNode1D.cc
hipStreamIsCapturing.cc
hipStreamGetCaptureInfo.cc
hipStreamEndCapture.cc
)
hip_add_exe_to_target(NAME GraphsTest
tests/catch/unit/graph/hipGraphAddHostNode.cc
+6 -4
Dosyayı Görüntüle
@@ -56,14 +56,16 @@ static void __global__ vector_square(int *B_d, int *D_d) {
}
}
static void vectorsquare_callback(void* ptr) {
// The callback func is not working with zero parameters
// Temporary fix for adding the below 2 lines and ticket
// has been raised for the same.
// The callback func is hipHostFn_t which is
// of type void (*)(void*). This test is designed to
// work with global variables, hence the workaround to
// print this *ptr value to avoid type mismatch errors.
int *A = reinterpret_cast<int *>(ptr);
A++;
for (int i = 0; i < SIZE; i++) {
if (D_h[i] != B_h[i] * B_h[i]) {
INFO("Validation failed " << D_h[i] << B_h[i]);
INFO("Ignore this garbage value" << *A);
REQUIRE(false);
}
}
tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc
+2
Dosyayı Görüntüle
@@ -273,6 +273,7 @@ void hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(bool device_ctxchg = false,
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
@@ -419,6 +420,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel") {
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc
+2
Dosyayı Görüntüle
@@ -250,6 +250,7 @@ void hipGraphAddMemcpyNodeToSymbol_GlobalMemory(bool device_ctxchg = false,
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
@@ -384,6 +385,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel") {
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
tests/catch/unit/graph/hipGraphClone.cc
+1
Dosyayı Görüntüle
@@ -295,6 +295,7 @@ TEST_CASE("Unit_hipGraphClone_MultiThreaded") {
HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
for (size_t i = 0; i < N; i++) {
if (A_h[i] != B_h[i]) {
tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc
+317
Dosyayı Görüntüle
@@ -0,0 +1,317 @@
/*
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 :
1) Set a different type of event using hipGraphEventWaitNodeSetEvent and
validate using hipGraphEventWaitNodeGetEvent.
2) Create a graph1 with memset (Value 1) node, event record node (event A)
and memset (Value 2) node and event record node (event B). Create a
graph2 with Event Wait (event A ) node and memcpyd2h. Instantiate
graph1 on stream1 and graph2 on stream2. Set the Event Wait node event
to B using hipGraphEventWaitNodeSetEvent. Launch graphs. Wait for the
event to complete. Verify the results.
3) Negative Scenarios
- Input node parameter is passed as nullptr.
- Input event parameter is passed as nullptr.
- Input node is an empty node.
- Input node is a memset node.
- Input node is a event record node.
- Input node is an uninitialized node.
- Input event is an uninitialized node.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define LEN 512
/**
* Local Function
*/
static void validateEventWaitNodeSetEvent(unsigned flag) {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create events
hipEvent_t event1, event2, event_out;
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreateWithFlags(&event2, flag));
hipGraphNode_t eventwait;
HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
event1));
// Set a different event
HIP_CHECK(hipGraphEventWaitNodeSetEvent(eventwait, event2));
HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out));
// validate set event and get event are same
REQUIRE(event2 == event_out);
// Free resources
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event1));
HIP_CHECK(hipEventDestroy(event2));
}
/**
* Local Function
*/
static void setEventRecordNode() {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create events
hipEvent_t event1, event2;
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
event1));
// Set a different event eventrec using hipGraphEventWaitNodeSetEvent
REQUIRE(hipErrorInvalidValue ==
hipGraphEventWaitNodeSetEvent(eventrec, event2));
// Free resources
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event1));
HIP_CHECK(hipEventDestroy(event2));
}
/**
* Scenario 2
*/
TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") {
size_t memsize = LEN * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, LEN);
size_t NElem{LEN};
hipGraph_t graph1, graph2;
hipStream_t streamForGraph1, streamForGraph2;
hipGraphExec_t graphExec1, graphExec2;
HIP_CHECK(hipStreamCreate(&streamForGraph1));
HIP_CHECK(hipGraphCreate(&graph1, 0));
HIP_CHECK(hipGraphCreate(&graph2, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph2));
hipEvent_t event1, event2;
HIP_CHECK(hipEventCreateWithFlags(&event1, hipEventDefault));
HIP_CHECK(hipEventCreateWithFlags(&event2, hipEventBlockingSync));
hipGraphNode_t event_rec_node, event_wait_node;
int *inp_h, *inp_d, *out_h_g1, *out_d_g1, *out_h_g2, *out_d_g2;
// Allocate host buffers
inp_h = reinterpret_cast<int*>(malloc(memsize));
REQUIRE(inp_h != nullptr);
out_h_g1 = reinterpret_cast<int*>(malloc(memsize));
REQUIRE(out_h_g1 != nullptr);
out_h_g2 = reinterpret_cast<int*>(malloc(memsize));
REQUIRE(out_h_g2 != nullptr);
// Allocate device buffers
HIP_CHECK(hipMalloc(&inp_d, memsize));
HIP_CHECK(hipMalloc(&out_d_g1, memsize));
HIP_CHECK(hipMalloc(&out_d_g2, memsize));
// Initialize host buffer
for (uint32_t i = 0; i < LEN; i++) {
inp_h[i] = i;
out_h_g1[i] = 0;
out_h_g2[i] = 0;
}
// Graph1 creation ...........
// Create event1 record node in graph1
HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0,
event1));
// Create memcpy and kernel nodes for graph1
hipGraphNode_t memcpyH2D, memcpyD2H_1, kernelnode_1;
hipKernelNodeParams kernelNodeParams1{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d,
inp_h, memsize, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0,
out_h_g1, out_d_g1, memsize, hipMemcpyDeviceToHost));
void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast<void *>(&NElem)};
kernelNodeParams1.func =
reinterpret_cast<void *>(HipTest::vector_square<int>);
kernelNodeParams1.gridDim = dim3(blocks);
kernelNodeParams1.blockDim = dim3(threadsPerBlock);
kernelNodeParams1.sharedMemBytes = 0;
kernelNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams1.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0,
&kernelNodeParams1));
// Create dependencies for graph1
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D,
&event_rec_node, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node,
&kernelnode_1, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1,
&memcpyD2H_1, 1));
// Graph2 creation ...........
// Create event1 record node in graph2
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0,
event1));
// Create memcpy and kernel nodes for graph2
hipGraphNode_t memcpyD2H_2, kernelnode_2;
hipKernelNodeParams kernelNodeParams2{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0,
out_h_g2, out_d_g2, memsize, hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast<void *>(&NElem)};
kernelNodeParams2.func =
reinterpret_cast<void *>(HipTest::vector_cubic<int>);
kernelNodeParams2.gridDim = dim3(blocks);
kernelNodeParams2.blockDim = dim3(threadsPerBlock);
kernelNodeParams2.sharedMemBytes = 0;
kernelNodeParams2.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams2.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0,
&kernelNodeParams2));
// Create dependencies for graph2
HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node,
&kernelnode_2, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2,
&memcpyD2H_2, 1));
// Instantiate and launch the graphs
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0));
// Set event
HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_rec_node, event2));
HIP_CHECK(hipGraphEventWaitNodeSetEvent(event_wait_node, event2));
HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1));
HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2));
HIP_CHECK(hipStreamSynchronize(streamForGraph1));
HIP_CHECK(hipStreamSynchronize(streamForGraph2));
// Validate output
bool btestPassed1 = true;
for (uint32_t i = 0; i < LEN; i++) {
if (out_h_g1[i] != (inp_h[i]*inp_h[i])) {
btestPassed1 = false;
break;
}
}
REQUIRE(btestPassed1 == true);
bool btestPassed2 = true;
for (uint32_t i = 0; i < LEN; i++) {
if (out_h_g2[i] != (inp_h[i]*inp_h[i]*inp_h[i])) {
btestPassed2 = false;
break;
}
}
REQUIRE(btestPassed2 == true);
// Destroy all resources
HIP_CHECK(hipFree(inp_d));
HIP_CHECK(hipFree(out_d_g1));
HIP_CHECK(hipFree(out_d_g2));
free(inp_h);
free(out_h_g1);
free(out_h_g2);
HIP_CHECK(hipGraphExecDestroy(graphExec1));
HIP_CHECK(hipGraphExecDestroy(graphExec2));
HIP_CHECK(hipGraphDestroy(graph1));
HIP_CHECK(hipGraphDestroy(graph2));
HIP_CHECK(hipEventDestroy(event1));
HIP_CHECK(hipEventDestroy(event2));
HIP_CHECK(hipStreamDestroy(streamForGraph1));
HIP_CHECK(hipStreamDestroy(streamForGraph2));
}
/**
* Scenario 1
*/
TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetGet") {
SECTION("Flag = hipEventDefault") {
validateEventWaitNodeSetEvent(hipEventDefault);
}
SECTION("Flag = hipEventBlockingSync") {
validateEventWaitNodeSetEvent(hipEventBlockingSync);
}
SECTION("Flag = hipEventDisableTiming") {
validateEventWaitNodeSetEvent(hipEventDisableTiming);
}
}
/**
* Scenario 3
*/
TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event1, event2;
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventwait;
HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
event1));
SECTION("node = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent(
nullptr, event2));
}
SECTION("event = nullptr") {
REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent(
eventwait, nullptr));
}
SECTION("input node is empty node") {
hipGraphNode_t EmptyGraphNode;
HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0));
REQUIRE(hipErrorInvalidValue ==
hipGraphEventWaitNodeSetEvent(EmptyGraphNode, event2));
}
SECTION("input node is memset node") {
constexpr size_t Nbytes = 1024;
char *A_d;
hipGraphNode_t memset_A;
hipMemsetParams memsetParams{};
HIP_CHECK(hipMalloc(&A_d, Nbytes));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
REQUIRE(hipErrorInvalidValue ==
hipGraphEventWaitNodeSetEvent(memset_A, event2));
HIP_CHECK(hipFree(A_d));
}
SECTION("input node is event record node") {
setEventRecordNode();
}
SECTION("input node is uninitialized node") {
hipGraphNode_t node_uninit{};
REQUIRE(hipErrorInvalidValue ==
hipGraphEventWaitNodeSetEvent(node_uninit, event2));
}
SECTION("input event is uninitialized") {
hipEvent_t event_uninit{};
REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent(
eventwait, event_uninit));
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event1));
HIP_CHECK(hipEventDestroy(event2));
}
tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc
+1 -1
Dosyayı Görüntüle
@@ -57,7 +57,7 @@ Testcase Scenarios :
* Kernel Functions to copy.
*/
static __global__ void copy_ker_func(int* a, int* b) {
int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
int tx = blockIdx.x*blockDim.x + threadIdx.x;
if (tx < LEN) b[tx] = a[tx];
}
tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc
+1
Dosyayı Görüntüle
@@ -269,6 +269,7 @@ void hipGraphExecMemcpyNodeSetParamsFromSymbol_GlobalMem(bool useConstVar) {
}
HIP_CHECK(hipGraphLaunch(graphExec, 0));
HIP_CHECK(hipStreamSynchronize(0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
tests/catch/unit/graph/hipStreamEndCapture.cc
+176
Dosyayı Görüntüle
@@ -0,0 +1,176 @@
/*
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.
*/
/**
Negative Testcase Scenarios :
1) Pass stream as nullptr and verify there is no crash, api returns error code.
2) Pass graph as nullptr and verify there is no crash, api returns error code.
3) Pass graph as nullptr and and stream as hipStreamPerThread verify there
is no crash, api returns error code.
4) End capture on stream where capture has not yet started and verify
error code is returned.
5) Destroy stream and try to end capture.
6) Destroy Graph and try to end capture.
7) Begin capture on a thread with mode other than hipStreamCaptureModeRelaxed
and try to end capture from different thread. Expect to return
hipErrorStreamCaptureWrongThread.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
TEST_CASE("Unit_hipStreamEndCapture_Negative") {
hipError_t ret;
SECTION("Pass stream as nullptr") {
hipGraph_t graph;
ret = hipStreamEndCapture(nullptr, &graph);
REQUIRE(hipErrorIllegalState == ret);
}
#if HT_NVIDIA
SECTION("Pass graph as nullptr") {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
ret = hipStreamEndCapture(stream, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
HIP_CHECK(hipStreamDestroy(stream));
}
SECTION("Pass graph as nullptr and stream as hipStreamPerThread") {
ret = hipStreamEndCapture(hipStreamPerThread, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
SECTION("End capture on stream where capture has not yet started") {
hipStream_t stream;
hipGraph_t graph;
HIP_CHECK(hipStreamCreate(&stream));
ret = hipStreamEndCapture(stream, &graph);
REQUIRE(hipErrorIllegalState == ret);
HIP_CHECK(hipStreamDestroy(stream));
}
SECTION("Destroy stream and try to end capture") {
hipStream_t stream;
hipGraph_t graph;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamDestroy(stream));
ret = hipStreamEndCapture(stream, &graph);
REQUIRE(hipErrorContextIsDestroyed == ret);
}
SECTION("Destroy graph and try to end capture in between") {
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
constexpr size_t N = 100000;
size_t Nbytes = N * sizeof(float);
float *A_d, *C_d;
float *A_h, *C_h;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipGraphDestroy(graph));
ret = hipStreamEndCapture(stream, &graph);
REQUIRE(hipSuccess == ret);
free(A_h);
free(C_h);
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipStreamDestroy(stream));
}
}
static void thread_func(hipStream_t stream, hipGraph_t graph) {
HIP_ASSERT(hipErrorStreamCaptureWrongThread ==
hipStreamEndCapture(stream, &graph));
}
TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") {
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
constexpr size_t N = 100000;
size_t Nbytes = N * sizeof(float);
float *A_d, *C_d;
float *A_h, *C_h;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
std::thread t(thread_func, stream, graph);
t.join();
#if HT_AMD
HIP_CHECK(hipStreamEndCapture(stream, &graph));
#endif
free(A_h);
free(C_h);
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipGraphDestroy(graph));
}
tests/catch/unit/graph/hipStreamGetCaptureInfo.cc
+231
Dosyayı Görüntüle
@@ -0,0 +1,231 @@
/*
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) Start stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusActive
and identifier returned is valid/non-zero.
2) End stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone
and identifier is not returned/updated by api.
3) Begin capture on hipStreamPerThread and get capture info. Verify api is success, capture status is hipStreamCaptureStatusActive
and identifier returned is valid/non-zero.
4) End capture on hipStreamPerThread, get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone
and identifier is not returned/updated by api.
5) Perform multiple captures and verify the identifier returned is unique.
Argument Validation/Negative:
1) Pass pId as nullptr and verify api doesn’t crash and returns success.
2) Pass pCaptureStatus as nullptr and verify api doesn’t crash and returns error code.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
constexpr size_t N = 1000000;
constexpr int LAUNCH_ITERS = 1;
/**
* Validates stream capture info, launches graph and verify results
*/
void validateStreamCaptureInfo(hipStream_t mstream) {
hipStream_t stream1{nullptr}, stream2{nullptr}, streamForLaunch{nullptr};
hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent;
hipGraph_t graph{nullptr};
hipGraphExec_t graphExec{nullptr};
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
size_t Nbytes = N * sizeof(float);
float *A_d, *C_d;
float *A_h, *C_h;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
HIP_CHECK(hipStreamCreate(&streamForLaunch));
// Initialize input buffer
for (size_t i = 0; i < N; ++i) {
A_h[i] = 3.146f + i; // Pi
}
// Create cross stream dependencies.
// memset operations are done on stream1 and stream2
// and they are joined back to mainstream
HIP_CHECK(hipStreamCreate(&stream1));
HIP_CHECK(hipStreamCreate(&stream2));
HIP_CHECK(hipEventCreate(&memsetEvent1));
HIP_CHECK(hipEventCreate(&memsetEvent2));
HIP_CHECK(hipEventCreate(&forkStreamEvent));
HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(forkStreamEvent, mstream));
HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0));
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
HIP_CHECK(hipEventRecord(memsetEvent1, stream1));
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2));
HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0));
HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, mstream, A_d, C_d, N);
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
unsigned long long capSequenceID = 0; // NOLINT
HIP_CHECK(hipStreamGetCaptureInfo(mstream, &captureStatus, &capSequenceID));
// verify capture status is active and sequence id is valid
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
REQUIRE(capSequenceID > 0);
// End capture and verify graph is returned
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream));
HIP_CHECK(hipStreamEndCapture(mstream, &graph));
REQUIRE(graph != nullptr);
// verify capture status is inactive and sequence id is not updated
capSequenceID = 0;
HIP_CHECK(hipStreamGetCaptureInfo(mstream, &captureStatus, &capSequenceID));
REQUIRE(captureStatus == hipStreamCaptureStatusNone);
REQUIRE(capSequenceID == 0);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < LAUNCH_ITERS; i++) {
HIP_CHECK(hipGraphLaunch(graphExec, streamForLaunch));
}
HIP_CHECK(hipStreamSynchronize(streamForLaunch));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForLaunch));
HIP_CHECK(hipStreamDestroy(stream1));
HIP_CHECK(hipStreamDestroy(stream2));
HIP_CHECK(hipEventDestroy(forkStreamEvent));
HIP_CHECK(hipEventDestroy(memsetEvent1));
HIP_CHECK(hipEventDestroy(memsetEvent2));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
// Validate the computation
for (size_t i = 0; i < N; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
INFO("A and C not matching at " << i << " C_h[i] " << C_h[i]
<< " A_h[i] " << A_h[i]);
REQUIRE(false);
}
}
free(A_h);
free(C_h);
}
/**
* Basic Functional Test for stream capture and getting capture info.
* Regular/custom stream is used for stream capture.
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_BasicFunctional") {
hipStream_t streamForCapture;
HIP_CHECK(hipStreamCreate(&streamForCapture));
validateStreamCaptureInfo(streamForCapture);
HIP_CHECK(hipStreamDestroy(streamForCapture));
}
/**
* Test performs stream capture on hipStreamPerThread and validates
* capture info.
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_hipStreamPerThread") {
validateStreamCaptureInfo(hipStreamPerThread);
}
/**
* Test starts stream capture on multiple streams and verifies uniqueness of
* identifiers returned.
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_UniqueID") {
constexpr int numStreams = 100;
hipStream_t streams[numStreams]{};
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
std::vector<int> idlist;
unsigned long long capSequenceID{}; // NOLINT
hipGraph_t graph{nullptr};
for (int i = 0; i < numStreams; i++) {
HIP_CHECK(hipStreamCreate(&streams[i]));
HIP_CHECK(hipStreamBeginCapture(streams[i], hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamGetCaptureInfo(streams[i], &captureStatus,
&capSequenceID));
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
REQUIRE(capSequenceID > 0);
idlist.push_back(capSequenceID);
}
for (int i = 0; i < numStreams; i++) {
for (int j = i+1; j < numStreams; j++) {
if (idlist[i] == idlist[j]) {
INFO("Same identifier returned for stream "
<< i << " and stream " << j);
REQUIRE(false);
}
}
}
for (int i = 0; i < numStreams; i++) {
HIP_CHECK(hipStreamEndCapture(streams[i], &graph));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streams[i]));
}
}
/**
* Argument validation/Negative tests for api
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_ArgValidation") {
hipError_t ret;
hipStream_t stream;
hipStreamCaptureStatus captureStatus;
unsigned long long capSequenceID; // NOLINT
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Capture ID location as nullptr") {
ret = hipStreamGetCaptureInfo(stream, &captureStatus, nullptr);
// Capture ID is optional
REQUIRE(ret == hipSuccess);
}
SECTION("Capture Status location as nullptr") {
ret = hipStreamGetCaptureInfo(stream, nullptr, &capSequenceID);
REQUIRE(ret == hipErrorInvalidValue);
}
HIP_CHECK(hipStreamDestroy(stream));
}
tests/catch/unit/graph/hipStreamIsCapturing.cc
+215
Dosyayı Görüntüle
@@ -0,0 +1,215 @@
/*
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_kernels.hh>
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
constexpr size_t N = 100000;
constexpr size_t Nbytes = N * sizeof(float);
/**
API - hipStreamIsCapturing
Negative Testcase Scenarios : Negative
1) Check capture status with null pCaptureStatus.
2) Check capture status with hipStreamPerThread and null pCaptureStatus.
Functional Testcase Scenarios :
1) Check capture status with null stream.
2) Check capture status with hipStreamPerThread.
3) Functional : Create a stream, call api and check
capture status is hipStreamCaptureStatusNone.
4) Functional : Start capturing a stream and check
capture status returned as hipStreamCaptureStatusActive.
5) Functional : Stop capturing a stream and check
status is returned as hipStreamCaptureStatusNone.
6) Functional : Use hipStreamPerThread, call api and check
capture status is hipStreamCaptureStatusNone.
7) Functional : Start capturing using hipStreamPerThread and check
capture status returned as hipStreamCaptureStatusActive.
8) Functional : Stop capturing using hipStreamPerThread and check
status is returned as hipStreamCaptureStatusNone.
*/
TEST_CASE("Unit_hipStreamIsCapturing_Negative") {
hipError_t ret;
hipStream_t stream{};
SECTION("Check capture status with null pCaptureStatus.") {
ret = hipStreamIsCapturing(stream, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Check capture status with hipStreamPerThread and"
" nullptr as pCaptureStatus.") {
ret = hipStreamIsCapturing(hipStreamPerThread, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
}
TEST_CASE("Unit_hipStreamIsCapturing_Functional_Basic") {
hipStreamCaptureStatus cStatus;
SECTION("Check capture status with null stream.") {
HIP_CHECK(hipStreamIsCapturing(nullptr, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
SECTION("Check capture status with hipStreamPerThread.") {
HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
}
/**
Testcase Scenarios :
1) Functional : Create a stream, call api and check
capture status is hipStreamCaptureStatusNone.
2) Functional : Start capturing a stream and check
capture status returned as hipStreamCaptureStatusActive.
3) Functional : Stop capturing a stream and check
status is returned as hipStreamCaptureStatusNone.
*/
TEST_CASE("Unit_hipStreamIsCapturing_Functional") {
float *A_d, *C_d;
float *A_h, *C_h;
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
hipStreamCaptureStatus cStatus;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Check the stream capture status before start capturing.") {
HIP_CHECK(hipStreamIsCapturing(stream, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
SECTION("Start capturing a stream and check the status.") {
HIP_CHECK(hipStreamIsCapturing(stream, &cStatus));
REQUIRE(hipStreamCaptureStatusActive == cStatus);
}
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamEndCapture(stream, &graph));
SECTION("Stop capturing a stream and check the status.") {
HIP_CHECK(hipStreamIsCapturing(stream, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
free(A_h);
free(C_h);
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
}
/**
Testcase Scenarios :
1) Functional : Use hipStreamPerThread, call api and check
capture status is hipStreamCaptureStatusNone.
2) Functional : Start capturing using hipStreamPerThread and check
capture status returned as hipStreamCaptureStatusActive.
3) Functional : Stop capturing using hipStreamPerThread and check
status is returned as hipStreamCaptureStatusNone.
*/
TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") {
float *A_d, *C_d;
float *A_h, *C_h;
hipGraph_t graph{nullptr};
hipStreamCaptureStatus cStatus;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
}
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
SECTION("Check the stream capture status before start capturing.") {
HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
HIP_CHECK(hipStreamBeginCapture(hipStreamPerThread,
hipStreamCaptureModeGlobal));
SECTION("Start capturing a stream and check the status.") {
HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus));
REQUIRE(hipStreamCaptureStatusActive == cStatus);
}
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
hipStreamPerThread));
HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, hipStreamPerThread));
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, hipStreamPerThread, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
hipStreamPerThread));
HIP_CHECK(hipStreamEndCapture(hipStreamPerThread, &graph));
SECTION("Stop capturing a stream and check the status.") {
HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus));
REQUIRE(hipStreamCaptureStatusNone == cStatus);
}
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
HIP_CHECK(hipGraphDestroy(graph));
free(A_h);
free(C_h);
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
}
tests/catch/unit/memory/CMakeLists.txt
+3 -6
Dosyayı Görüntüle
@@ -41,6 +41,7 @@ set(TEST_SRC
hipMemPtrGetInfo.cc
hipPointerGetAttributes.cc
hipHostGetFlags.cc
hipHostGetDevicePointer.cc
hipMallocManaged_MultiScenario.cc
hipMemsetNegative.cc
hipMemset.cc
@@ -57,10 +58,8 @@ set(TEST_SRC
hipMallocManaged.cc
hipMemRangeGetAttribute.cc
hipMemcpyFromSymbol.cc
hipMemcpyFromSymbolAsync.cc
hipMemcpyToSymbol.cc
hipMemcpyToSymbolAsync.cc
hipPtrGetAttribute.cc
hipMemPoolApi.cc
hipMemcpyPeer.cc
hipMemcpyPeerAsync.cc
hipMemcpyWithStream.cc
@@ -106,6 +105,7 @@ set(TEST_SRC
hipMemcpy_MultiThread.cc
hipHostRegister.cc
hipHostGetFlags.cc
hipHostGetDevicePointer.cc
hipMallocManaged_MultiScenario.cc
hipMemsetNegative.cc
hipMemset.cc
@@ -121,9 +121,6 @@ set(TEST_SRC
hipMallocManaged.cc
hipMemRangeGetAttribute.cc
hipMemcpyFromSymbol.cc
hipMemcpyFromSymbolAsync.cc
hipMemcpyToSymbol.cc
hipMemcpyToSymbolAsync.cc
hipPtrGetAttribute.cc
hipMemcpyPeer.cc
hipMemcpyPeerAsync.cc
tests/catch/unit/memory/hipHostGetDevicePointer.cc
+78
Dosyayı Görüntüle
@@ -0,0 +1,78 @@
/*
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>
TEST_CASE("Unit_hipHostGetDevicePointer_Negative") {
int* hPtr{nullptr};
HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));
SECTION("Nullptr as device") {
HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue);
}
SECTION("Nullptr as host") {
int* dPtr{nullptr};
HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), nullptr, 0),
hipErrorInvalidValue);
}
// Not adding check for flags since CUDA spec states that there might be more values added to it
HIP_CHECK(hipHostFree(hPtr));
}
template <typename T> __global__ void set(T* ptr, T val) { *ptr = val; }
TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") {
int* hPtr{nullptr};
HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));
auto kernel = set<int>;
constexpr int value = 10;
SECTION("Set the value on device - Get device ptr") {
int* dPtr{nullptr};
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 0));
REQUIRE(dPtr != nullptr);
kernel<<<1, 1>>>(dPtr, value);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(*hPtr == value);
}
SECTION("Set the value on device - by hipHostRegister") {
int res{0}; // Stuff on stack
HIP_CHECK(hipHostRegister(&res, sizeof(int), 0)); // Lets map stack memory :)
int* dPtr{nullptr};
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), &res, 0))
kernel<<<1, 1>>>(dPtr, value);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipHostUnregister(&res));
REQUIRE(value == 10);
}
HIP_CHECK(hipHostFree(hPtr));
}
tests/catch/unit/memory/hipHostMalloc.cc
+20 -10
Dosyayı Görüntüle
@@ -30,8 +30,9 @@ This testfile verifies the following scenarios of hipHostMalloc API
*/
#include<hip_test_checkers.hh>
#include<hip_test_kernels.hh>
#include<kernels.hh>
#include<hip_test_common.hh>
#include <hip_test_context.hh>
#define SYNC_EVENT 0
#define SYNC_STREAM 1
@@ -41,11 +42,6 @@ std::vector<std::string> syncMsg = {"event", "stream", "device"};
static constexpr int numElements{1024 * 16};
static constexpr size_t sizeBytes{numElements * sizeof(int)};
__global__ void Set(int* Ad, int val) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
Ad[tx] = val;
}
void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags,
int syncMethod, std::string msg) {
std::cerr << "test: CheckHostPointer "
@@ -70,10 +66,10 @@ void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags,
const int expected = 13;
// Init array to know state:
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42);
HipTest::launchKernel(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42);
HIP_CHECK(hipDeviceSynchronize());
hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, ptr, expected);
HipTest::launchKernel(Set, dimGrid, dimBlock, 0, s, ptr, expected);
HIP_CHECK(hipEventRecord(e, s));
// Host waits for event :
@@ -137,9 +133,9 @@ TEST_CASE("Unit_hipHostMalloc_Basic") {
dim3 dimGrid(LEN / 512, 1, 1);
dim3 dimBlock(512, 1, 1);
hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock,
HipTest::launchKernel<float>(HipTest::vectorADD<float>, dimGrid, dimBlock,
0, 0, static_cast<const float*>(A_d),
static_cast<const float*>(B_d), C_d, LEN);
static_cast<const float*>(B_d), C_d, static_cast<size_t>(LEN));
HIP_CHECK(hipMemcpy(C_h, C_d, LEN*sizeof(float),
hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
@@ -148,6 +144,7 @@ TEST_CASE("Unit_hipHostMalloc_Basic") {
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(B_h));
HIP_CHECK(hipHostFree(C_h));
TestContext::get().cleanContext();
}
}
/*
@@ -185,6 +182,7 @@ TEST_CASE("Unit_hipHostMalloc_NonCoherent") {
SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem,
SYNC_EVENT, ptrType);
TestContext::get().cleanContext();
}
/*
@@ -233,4 +231,16 @@ TEST_CASE("Unit_hipHostMalloc_Default") {
CheckHostPointer(numElements, A, 0, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, 0, SYNC_EVENT, ptrType);
}
TEST_CASE("Unit_hipHostGetDevicePointer_NullCheck") {
int* d_a;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&d_a), sizeof(int)));
auto res = hipHostGetDevicePointer(nullptr,d_a,0);
REQUIRE(res == hipErrorInvalidValue);
HIP_CHECK(hipHostFree(d_a));
}
tests/catch/unit/memory/hipHostMallocTests.cc
+17 -10
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -29,33 +29,40 @@ Testcase Scenarios :
*/
#include <hip_test_common.hh>
#include <hip_test_helper.hh>
/**
* Performs argument validation of hipHostMalloc api.
*/
TEST_CASE("Unit_hipHostMalloc_ArgValidation") {
hipError_t ret;
#if HT_NVIDIA
HipTest::HIP_SKIP_TEST("TODO: Need to debug");
#endif
constexpr size_t allocSize = 1000;
char *ptr;
char* ptr;
SECTION("Pass ptr as nullptr") {
ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize);
REQUIRE(ret != hipSuccess);
HIP_CHECK_ERROR(hipHostMalloc(static_cast<void**>(nullptr), allocSize), hipErrorInvalidValue);
}
SECTION("Size as max(size_t)") {
ret = hipHostMalloc(&ptr, std::numeric_limits<std::size_t>::max());
REQUIRE(ret != hipSuccess);
HIP_CHECK_ERROR(hipHostMalloc(&ptr, (std::numeric_limits<std::size_t>::max)()),
hipErrorMemoryAllocation);
}
SECTION("Flags as max(uint)") {
ret = hipHostMalloc(&ptr, allocSize,
std::numeric_limits<unsigned int>::max());
REQUIRE(ret != hipSuccess);
HIP_CHECK_ERROR(hipHostMalloc(&ptr, allocSize, (std::numeric_limits<unsigned int>::max)()),
hipErrorInvalidValue);
}
SECTION("Pass size as zero and check ptr reset") {
HIP_CHECK(hipHostMalloc(&ptr, 0));
REQUIRE(ptr == nullptr);
}
SECTION("Pass hipHostMallocCoherent and hipHostMallocNonCoherent simultaneously") {
HIP_CHECK_ERROR(
hipHostMalloc(&ptr, allocSize, hipHostMallocCoherent | hipHostMallocNonCoherent),
hipErrorInvalidValue);
}
}
tests/catch/unit/memory/hipMemCoherencyTst.cc
+29 -37
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -34,17 +34,9 @@
#include <hip_test_common.hh>
#include <chrono>
__global__ void CoherentTst(int *ptr, int PeakClk) {
// Incrementing the value by 1
int64_t GpuFrq = (PeakClk * 1000);
int64_t StrtTck = clock64();
atomicAdd(ptr, 1);
// The following while loop checks the value in ptr for around 3-4 seconds
while ((clock64() - StrtTck) <= (3 * GpuFrq)) {
if (*ptr == 3) {
atomicAdd(ptr, 1);
return;
}
__global__ void CoherentTst(int* ptr) { // ptr was set to 1
atomicAdd(ptr, 1); // now ptr is 2
while (atomicCAS(ptr, 3, 4) != 3) { // wait till ptr is 3, then change it to 4
}
}
@@ -59,34 +51,34 @@ __global__ void SquareKrnl(int *ptr) {
static bool YES_COHERENT = false;
// The function tests the coherency of allocated memory
static void TstCoherency(int *Ptr, bool HmmMem) {
int *Dptr = nullptr, peak_clk;
hipStream_t strm;
HIP_CHECK(hipStreamCreate(&strm));
// If this test hangs, means there is issue in coherency
static void TstCoherency(int* ptr, bool hmmMem) {
int* dptr = nullptr;
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
// storing value 1 in the memory created above
*Ptr = 1;
// Getting gpu frequency
HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
if (!HmmMem) {
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr), Ptr,
0));
CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk);
*ptr = 1;
if (!hmmMem) {
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&dptr), ptr, 0));
CoherentTst<<<1, 1, 0, stream>>>(dptr);
} else {
CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk);
CoherentTst<<<1, 1, 0, stream>>>(ptr);
}
// looping until the value is 2 for 3 seconds
std::chrono::steady_clock::time_point start =
std::chrono::steady_clock::now();
while (std::chrono::duration_cast<std::chrono::seconds>(
std::chrono::steady_clock::now() - start).count() < 3) {
if (*Ptr == 2) {
*Ptr += 1;
break;
}
}
HIP_CHECK(hipStreamSynchronize(strm));
HIP_CHECK(hipStreamDestroy(strm));
if (*Ptr == 4) {
std::chrono::steady_clock::time_point start = std::chrono::steady_clock::now();
while (std::chrono::duration_cast<std::chrono::seconds>(std::chrono::steady_clock::now() - start)
.count() < 3 &&
*ptr == 2) {
} // wait till ptr is 2 from kernel or 3 seconds
*ptr += 1; // increment it to 3
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
if (*ptr == 4) {
YES_COHERENT = true;
}
}
tests/catch/unit/memory/hipMemPoolApi.cc
+545
Dosyayı Görüntüle
@@ -0,0 +1,545 @@
/*
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, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* Test Case Description:
1) This testcase verifies the basic scenario - supported on
all devices
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <cstdio>
#include <cstdint>
#include <algorithm>
#include <thread>
#include <chrono>
constexpr hipMemPoolProps kPoolProps = {
hipMemAllocationTypePinned,
hipMemHandleTypeNone,
{
hipMemLocationTypeDevice,
0
},
nullptr,
{0}
};
/*
This testcase verifies HIP Mem Pool API basic scenario - supported on all devices
*/
TEST_CASE("Unit_hipMemPoolApi_Basic") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
int numElements = 64 * 1024 * 1024;
float *A = nullptr, *B = nullptr;
hipMemPool_t mem_pool = nullptr;
int device = 0;
HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, device));
HIP_CHECK(hipDeviceSetMemPool(device, mem_pool));
HIP_CHECK(hipDeviceGetMemPool(&mem_pool, device));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMallocAsync(&A, numElements * sizeof(float), stream));
INFO("hipMallocAsync result: " << A);
HIP_CHECK(hipFreeAsync(A, stream));
// Reset the default memory pool usage for the following tests
hipMemPoolAttr attr = hipMemPoolAttrUsedMemHigh;
std::uint64_t value64 = 0;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value64));
size_t min_bytes_to_hold = 1024 * 1024;
HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold));
attr = hipMemPoolReuseFollowEventDependencies;
int value = 0;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value));
hipMemAccessDesc desc_list = {};
int count = 1;
HIP_CHECK(hipMemPoolSetAccess(mem_pool, &desc_list, count));
hipMemAccessFlags flags = hipMemAccessFlagsProtNone;
hipMemLocation location = {};
HIP_CHECK(hipMemPoolGetAccess(&flags, mem_pool, &location));
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
HIP_CHECK(hipMallocFromPoolAsync(&B, numElements * sizeof(float), mem_pool, stream));
HIP_CHECK(hipMemPoolDestroy(mem_pool));
HIP_CHECK(hipStreamDestroy(stream));
}
constexpr auto wait_ms = 500;
__global__ void kernel500ms(float* hostRes, int clkRate) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hostRes[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
uint64_t start = clock64()/clkRate, cur;
if (clkRate > 1) {
do { cur = clock64()/clkRate-start;}while (cur < wait_ms);
} else {
do { cur = clock64()/start;}while (cur < wait_ms);
}
}
TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
hipMemPool_t mem_pool;
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
float* B, *C;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
size_t numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream));
numElements = 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), mem_pool, stream));
int blocks = 1024;
int clkRate;
hipMemPoolAttr attr;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
attr = hipMemPoolAttrReservedMemCurrent;
std::uint64_t res_before_sync = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_before_sync));
HIP_CHECK(hipStreamSynchronize(stream));
std::uint64_t res_after_sync = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_sync));
// Sync must releaae memory to OS
REQUIRE(res_after_sync < res_before_sync);
int value = 0;
attr = hipMemPoolReuseFollowEventDependencies;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value));
// Default enabled
REQUIRE(1 == value);
attr = hipMemPoolReuseAllowOpportunistic;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value));
// Default enabled
REQUIRE(1 == value);
attr = hipMemPoolReuseAllowInternalDependencies;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value));
// Default enabled
REQUIRE(1 == value);
attr = hipMemPoolAttrReleaseThreshold;
std::uint64_t value64 = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Default is 0
REQUIRE(0 == value64);
attr = hipMemPoolAttrReservedMemHigh;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Must be bigger than current
REQUIRE(value64 > res_after_sync);
attr = hipMemPoolAttrUsedMemCurrent;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current usage query works - just small buffer left
REQUIRE(sizeof(float) * 1024 == value64);
attr = hipMemPoolAttrUsedMemHigh;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the high watermark usage works - the both buffers must be reported
REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
HIP_CHECK(hipMemPoolDestroy(mem_pool));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemPoolApi_BasicTrim") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
hipMemPool_t mem_pool;
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
float* B, *C;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
size_t numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream));
numElements = 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), mem_pool, stream));
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
hipMemPoolAttr attr;
attr = hipMemPoolAttrReleaseThreshold;
// The pool must hold 128MB
std::uint64_t threshold = 128 * 1024 * 1024;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &threshold));
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
// Get reserved memory before trim
attr = hipMemPoolAttrReservedMemCurrent;
std::uint64_t res_before_trim = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_before_trim));
size_t min_bytes_to_hold = sizeof(float) * 1024;
HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold));
std::uint64_t res_after_trim = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_trim));
// Trim must be a nop because execution isn't done
REQUIRE(res_before_trim == res_after_trim);
HIP_CHECK(hipStreamSynchronize(stream));
std::uint64_t res_after_sync = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_sync));
// Since hipMemPoolAttrReleaseThreshold is 128 MB sync does nothing to the freed memory
REQUIRE(res_after_trim == res_after_sync);
HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold));
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_trim));
// Validate memory after real trim. The pool must hold less memory than before
REQUIRE(res_after_trim < res_after_sync);
attr = hipMemPoolAttrReleaseThreshold;
std::uint64_t value64 = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the threshold query works
REQUIRE(threshold == value64);
attr = hipMemPoolAttrUsedMemCurrent;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current usage query works - just small buffer left
REQUIRE(sizeof(float) * 1024 == value64);
attr = hipMemPoolAttrUsedMemHigh;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the high watermark usage works - the both buffers must be reported
REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
HIP_CHECK(hipMemPoolDestroy(mem_pool));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemPoolApi_BasicReuse") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
hipMemPool_t mem_pool;
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
float *A, *B, *C;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
size_t numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
numElements = 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), mem_pool, stream));
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream));
// Runtime must reuse the pointer
REQUIRE(A == B);
// Make a sync before the second kernel launch to make sure memory B isn't gone
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
HIP_CHECK(hipStreamSynchronize(stream));
attr = hipMemPoolAttrUsedMemCurrent;
std::uint64_t value64 = 0;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current usage reports the both buffers
REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
attr = hipMemPoolAttrUsedMemHigh;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the high watermark usage works - the both buffers must be reported
REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
attr = hipMemPoolAttrUsedMemCurrent;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current usage reports just one buffer, because the above free doesn't hold memory
REQUIRE(sizeof(float) * 1024 == value64);
HIP_CHECK(hipMemPoolDestroy(mem_pool));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemPoolApi_Opportunistic") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
hipMemPool_t mem_pool;
HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps));
hipMemPoolAttr attr;
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
float *A, *B, *C;
hipStream_t stream, stream2;
// Create 2 async non-blocking streams
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamNonBlocking));
size_t numElements = 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), mem_pool, stream));
int value = 0;
SECTION("Disallow Opportunistic - No Reuse") {
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
// Disable all default pool states
attr = hipMemPoolReuseFollowEventDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
attr = hipMemPoolReuseAllowOpportunistic;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
attr = hipMemPoolReuseAllowInternalDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
// Sleep for 1 second GPU should be idle by now
std::this_thread::sleep_for(std::chrono::milliseconds(1000));
numElements = 8 * 1024 * 1024;
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream2));
// Without Opportunistic state runtime must allocate another buffer
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream2));
}
SECTION("Allow Opportunistic - Reuse") {
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
value = 1;
attr = hipMemPoolReuseAllowOpportunistic;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
// Sleep for 1 second GPU should be idle by now
std::this_thread::sleep_for(std::chrono::milliseconds(1000));
numElements = 8 * 1024 * 1024;
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream2));
// With Opportunistic state runtime will reuse freed buffer A
REQUIRE(A == B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream2));
}
SECTION("Allow Opportunistic - No Reuse") {
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), mem_pool, stream));
value = 1;
attr = hipMemPoolReuseAllowOpportunistic;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value));
// Run kernel for 500 ms in the first stream
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
numElements = 8 * 1024 * 1024;
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), mem_pool, stream2));
// With Opportunistic state runtime can't reuse freed buffer A, because it's still busy with the kernel
REQUIRE(A != B);
// Run kernel with the new memory in the second stream
kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamSynchronize(stream2));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream2));
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipMemPoolDestroy(mem_pool));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipStreamDestroy(stream2));
}
TEST_CASE("Unit_hipMemPoolApi_Default") {
int mem_pool_support = 0;
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
hipMemPool_t mem_pool;
HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0));
float *A, *B, *C;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
size_t numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&A), numElements * sizeof(float), stream));
numElements = 1024;
HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&C), numElements * sizeof(float), stream));
int blocks = 2;
int clkRate;
HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate);
hipMemPoolAttr attr;
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(A), stream));
numElements = 8 * 1024 * 1024;
HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&B), numElements * sizeof(float), stream));
// Runtime must reuse the pointer
REQUIRE(A == B);
// Make a sync before the second kernel launch to make sure memory B isn't gone
HIP_CHECK(hipStreamSynchronize(stream));
// Second kernel launch with new memory
kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(B), stream));
HIP_CHECK(hipStreamSynchronize(stream));
std::uint64_t value64 = 0;
attr = hipMemPoolAttrReservedMemCurrent;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current reserved just 4KB alloc
REQUIRE(sizeof(float) * 1024 == value64);
attr = hipMemPoolAttrUsedMemHigh;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the high watermark usage works - the both buffers must be reported
REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64);
attr = hipMemPoolAttrUsedMemCurrent;
HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64));
// Make sure the current usage reports just one buffer, because the above free doesn't hold memory
REQUIRE(sizeof(float) * 1024 == value64);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(C), stream));
HIP_CHECK(hipStreamDestroy(stream));
}
tests/catch/unit/memory/hipMemPrefetchAsync.cc
+2 -2
Dosyayı Görüntüle
@@ -20,8 +20,8 @@ THE SOFTWARE.
#include <hip_test_common.hh>
// Kernel function
__global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) {
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < N; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc
+2 -2
Dosyayı Görüntüle
@@ -37,8 +37,8 @@ THE SOFTWARE.
// Kernel function
__global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) {
size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
size_t stride = hipBlockDim_x * hipGridDim_x;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = offset; i < N; i += stride) {
Hmm[i] = Hmm[i] * Hmm[i];
}
tests/catch/unit/memory/hipMemcpyFromSymbol.cc
+163 -16
Dosyayı Görüntüle
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -18,26 +18,173 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define SIZE 1024
/* Test verifies hipMemcpyFromSymbol API Negative scenarios.
__device__ int devSymbol[10];
/* Test verifies hipMemcpy[From/To]Symbol[Async] API Negative scenarios.
*/
TEST_CASE("Unit_hipMemcpyFromSymbol_Negative") {
void *Sd;
char S[SIZE]="This is not a device symbol";
HIP_CHECK(hipMalloc(&Sd, SIZE));
SECTION("Passing void pointer") {
REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd),
SIZE, 0, hipMemcpyDeviceToHost));
TEST_CASE("Unit_hipMemcpyFromToSymbol_Negative") {
SECTION("Invalid Src Ptr") {
int result{0};
HIP_CHECK_ERROR(
hipMemcpyFromSymbol(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyDeviceToHost),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyToSymbol(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice),
hipErrorInvalidSymbol);
HIP_CHECK_ERROR(
hipMemcpyToSymbolAsync(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice, nullptr),
hipErrorInvalidSymbol);
HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0,
hipMemcpyDeviceToHost, nullptr),
hipErrorInvalidValue);
}
SECTION("Passing NULL Pointer") {
REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, nullptr,
SIZE, 0, hipMemcpyDeviceToHost));
SECTION("Invalid Dst Ptr") {
int result{0};
HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost),
hipErrorInvalidSymbol);
HIP_CHECK_ERROR(
hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, hipMemcpyHostToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0,
hipMemcpyHostToDevice, nullptr),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpyFromSymbolAsync(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost, nullptr),
hipErrorInvalidSymbol);
}
HIP_CHECK(hipFree(Sd));
SECTION("Invalid Size") {
int result{0};
HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0,
hipMemcpyDeviceToHost),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0,
hipMemcpyHostToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0,
hipMemcpyHostToDevice, nullptr),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0,
hipMemcpyDeviceToHost, nullptr),
hipErrorInvalidValue);
}
SECTION("Invalid Offset") {
int result{0};
HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300,
hipMemcpyDeviceToHost),
hipErrorInvalidValue);
HIP_CHECK_ERROR(
hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, hipMemcpyHostToDevice),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300,
hipMemcpyHostToDevice, nullptr),
hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300,
hipMemcpyDeviceToHost, nullptr),
hipErrorInvalidValue);
}
SECTION("Invalid Direction") {
int result{0};
HIP_CHECK_ERROR(
hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyHostToDevice),
hipErrorInvalidMemcpyDirection);
HIP_CHECK_ERROR(
hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, hipMemcpyDeviceToHost),
hipErrorInvalidMemcpyDirection);
HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0,
hipMemcpyDeviceToHost, nullptr),
hipErrorInvalidMemcpyDirection);
HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0,
hipMemcpyHostToDevice, nullptr),
hipErrorInvalidMemcpyDirection);
}
}
/*
* Test Verifies hipMemcpyToSymbol/hipMemcpyFromSymbol and Async Variants for simple use case
* For single valuea To and From Symbol
* For Array Values To and From Symbol
* For Array Values with offset To and From Symbol
* For Sync and Async Variants*/
TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
enum StreamTestType { NullStream = 0, StreamPerThread, CreatedStream, NoStream };
/* Test type NoStream - Use Sync variants, else use async variants */
auto streamType = GENERATE(StreamTestType::NoStream, StreamTestType::NullStream,
StreamTestType::StreamPerThread, StreamTestType::CreatedStream);
hipStream_t stream{nullptr};
if (streamType == StreamTestType::StreamPerThread) {
stream = hipStreamPerThread;
} else if (streamType == StreamTestType::CreatedStream) {
HIP_CHECK(hipStreamCreate(&stream));
}
INFO("Stream :: " << streamType);
SECTION("Singular Value") {
int set{42};
int result{0};
if (streamType == StreamTestType::NoStream) {
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int)));
HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int)));
} else {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
REQUIRE(result == set);
}
SECTION("Array Values") {
constexpr size_t size{10};
int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2};
int result[size] = {0};
if (streamType == StreamTestType::NoStream) {
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size));
HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size));
} else {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
for (size_t i = 0; i < size; i++) {
REQUIRE(result[i] == set[i]);
}
}
SECTION("Offset'ed Values") {
constexpr size_t size{10};
constexpr size_t offset = 5 * sizeof(int);
int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2};
int result[size] = {0};
if (streamType == StreamTestType::NoStream) {
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset));
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset));
HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size));
} else {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, offset, 0, hipMemcpyHostToDevice,
stream));
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set + 5, offset, offset,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(result, HIP_SYMBOL(devSymbol), offset, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(result + 5, HIP_SYMBOL(devSymbol), offset, offset,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
}
for (size_t i = 0; i < size; i++) {
REQUIRE(result[i] == set[i]);
}
}
}
tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc
-47
Dosyayı Görüntüle
@@ -1,47 +0,0 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define SIZE 1024
/* Test verifies hipMemcpyFromSymbolAsync API Negative scenarios.
*/
TEST_CASE("Unit_hipMemcpyFromSymbolAsync_Negative") {
void *Sd;
char S[SIZE]="This is not a device symbol";
HIP_CHECK(hipMalloc(&Sd, SIZE));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Passing void pointer") {
REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd),
SIZE, 0, hipMemcpyDeviceToHost, stream));
}
SECTION("Passing NULL pointer") {
REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, nullptr,
SIZE, 0, hipMemcpyDeviceToHost, stream));
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(Sd));
}
tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc
-47
Dosyayı Görüntüle
@@ -1,47 +0,0 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define SIZE 1024
/* Test verifies hipMemcpyToSymbolAsync API Negative scenarios.
*/
TEST_CASE("Unit_hipMemcpyToSymbolAsync_Negative") {
void *Sd;
char S[SIZE]="This is not a device symbol";
HIP_CHECK(hipMalloc(&Sd, SIZE));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Passing void pointer") {
REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S,
SIZE, 0, hipMemcpyHostToDevice, stream));
}
SECTION("Passing NULL pointer") {
REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(nullptr, S,
SIZE, 0, hipMemcpyHostToDevice, stream));
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(Sd));
}
tests/catch/unit/stream/hipStreamGetPriority.cc
+44
Dosyayı Görüntüle
@@ -89,6 +89,50 @@ TEST_CASE("Unit_hipStreamGetPriority_happy") {
}
}
/**
* both stream and priority passed as nullptr.
*/
TEST_CASE("Unit_hipStreamGetPriority_nullptr_nullptr") {
auto res = hipStreamGetPriority(nullptr,nullptr);
REQUIRE(res == hipErrorInvalidValue);
}
/**
* valid stream and priority passed as nullptr.
*/
TEST_CASE("Unit_hipStreamGetPriority_stream_nullptr") {
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
auto res = hipStreamGetPriority(stream,nullptr);
REQUIRE(res == hipErrorInvalidValue);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* nullptr stream and valid priority
*/
TEST_CASE("Unit_hipStreamGetPriority_nullptr_priority") {
int priority = -1;
HIP_CHECK(hipStreamGetPriority(nullptr,&priority));
}
/**
* both stream and priority passed as valid.
*/
TEST_CASE("Unit_hipStreamGetPriority_stream_priority") {
int priority = -1;
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamGetPriority(stream,&priority));
HIP_CHECK(hipStreamDestroy(stream));
}
#if HT_AMD
/**
* Create stream with CUMask and check priority is returned as expected.
tests/catch/unit/stream/streamCommon.cc
+3 -3
Dosyayı Görüntüle
@@ -95,11 +95,11 @@ __global__ void waiting_kernel(int* semaphore) {
std::thread startSignalingThread(int* semaphore) {
std::thread signalingThread([semaphore]() {
hipStream_t signalingStream;
HIP_CHECK(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking));
HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking));
signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore);
HIP_CHECK(hipStreamSynchronize(signalingStream));
HIP_CHECK(hipStreamDestroy(signalingStream));
HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream));
HIP_CHECK_THREAD(hipStreamDestroy(signalingStream));
});
return signalingThread;
tests/catch/unit/stream/streamCommon.hh
+1
Dosyayı Görüntüle
@@ -44,6 +44,7 @@ __global__ void waiting_kernel(int* semaphore = nullptr);
/**
* @brief Creates a thread that runs a signaling_kernel on a non-blocking stream.
* hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream.
* You must call HIP_CHECK_THREAD_FINALIZE after joining this thread.
*
* @param semaphore memory location to signal
* @return std::thread thread that has to be joined after the testing is done.
tests/catch/unit/texture/hipBindTex2DPitch.cc
+2 -2
Dosyayı Görüntüle
@@ -28,8 +28,8 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;
// texture object is a kernel argument
static __global__ void texture2dCopyKernel(TYPE_t* dst) {
int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if ( (x < SIZE_W) && (y < SIZE_H) ) {
dst[SIZE_W*y+x] = tex2D(tex, x, y);
}
tests/catch/unit/texture/hipNormalizedFloatValueTex.cc
+1 -1
Dosyayı Görüntüle
@@ -43,7 +43,7 @@ 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;
unsigned int elementID = threadIdx.x;
if (elementID >= numElements)
return;
float coord = elementID/static_cast<float>(numElements);
tests/src/deviceLib/hipDeviceMalloc.cpp
+10 -10
Dosyayı Görüntüle
@@ -17,8 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM all
* TEST: %t EXCLUDE_HIP_PLATFORM all
* BUILD: %t %s ../../src/test_common.cpp
* TEST: %t
* HIT_END
*/
#include "test_common.h"
@@ -82,10 +82,10 @@ __device__ __host__ FloatT calc(FloatT A, FloatT B, enum CalcKind CK) {
// Copy value from A, B to allocated memory.
template <typename FloatT>
__global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) {
int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+ (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+ (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
* hipBlockDim_y;
int tx = threadIdx.x + blockDim.x * blockIdx.x
+ (threadIdx.y + blockDim.y * blockIdx.y) * blockDim.x
+ (threadIdx.z + blockDim.z * blockIdx.z) * blockDim.x
* blockDim.y;
if (tx == 0) {
*pA = (FloatT*)malloc(sizeof(FloatT) * LEN);
*pB = (FloatT*)malloc(sizeof(FloatT) * LEN);
@@ -100,10 +100,10 @@ __global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) {
// containing the address of the device-side allocated array.
template <typename FloatT>
__global__ void kernel_free(FloatT** pA, FloatT** pB, FloatT* C, enum CalcKind CK) {
int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x
+ (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x
+ (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x
* hipBlockDim_y;
int tx = threadIdx.x + blockDim.x * blockIdx.x
+ (threadIdx.y + blockDim.y * blockIdx.y) * blockDim.x
+ (threadIdx.z + blockDim.z * blockIdx.z) * blockDim.x
* blockDim.y;
C[tx] = calc<FloatT>((*pA)[tx], (*pB)[tx], CK);
if (tx == 0) {
free(*pA);
tests/src/runtimeApi/streamOperations/hipstream_operations.cpp
+30 -2
Dosyayı Görüntüle
@@ -58,6 +58,7 @@ constexpr unsigned int writeFlag = 0;
constexpr float SLEEP_MS = 100;
constexpr uint32_t DATA_INIT = 0x1234;
constexpr uint32_t DATA_UPDATE = 0X4321;
constexpr unsigned int ARR_SIZE = 5;
struct TEST_WAIT {
int compareOp;
@@ -203,31 +204,49 @@ void testWrite() {
uint64_t* host_ptr64 = (uint64_t *) malloc(sizeof(uint64_t));
uint32_t* host_ptr32 = (uint32_t *) malloc(sizeof(uint32_t));
uint64_t h_mem64[ARR_SIZE];
uint32_t h_mem32[ARR_SIZE];
uint64_t* h_mem64ptr = h_mem64;
uint32_t* h_mem32ptr = h_mem32;
std::cout << " hipStreamWriteValue: testing ... \n";
HIPCHECK(hipExtMallocWithFlags((void **)&signalPtr, 8, hipMallocSignalMemory));
void* device_ptr64;
void* device_ptr32;
uint64_t* d_mem64ptr;
uint32_t* d_mem32ptr;
*host_ptr64 = 0x0;
*host_ptr32 = 0x0;
*signalPtr = 0x0;
hipHostRegister(host_ptr64, sizeof(uint64_t), 0);
hipHostRegister(host_ptr32, sizeof(uint32_t), 0);
hipHostRegister(h_mem64ptr, sizeof(uint64_t) * ARR_SIZE, 0);
hipHostRegister(h_mem32ptr, sizeof(uint32_t) * ARR_SIZE, 0);
// Test writting registered host pointer
HIPCHECK(hipStreamWriteValue64(stream, host_ptr64, value64, writeFlag));
HIPCHECK(hipStreamWriteValue32(stream, host_ptr32, value32, writeFlag));
//test writting to an array
for (int indx = 0; indx < ARR_SIZE; indx++) {
HIPCHECK(hipStreamWriteValue64(stream, h_mem64ptr + indx, ARR_SIZE - indx, writeFlag));
HIPCHECK(hipStreamWriteValue32(stream, h_mem32ptr + indx, ARR_SIZE - indx, writeFlag));
}
hipStreamSynchronize(stream);
HIPASSERT(*host_ptr64 == value64);
HIPASSERT(*host_ptr32 == value32);
for (int indx = 0; indx < ARR_SIZE ; indx++) {
HIPASSERT(*(h_mem64ptr + indx) == (ARR_SIZE - indx));
HIPASSERT(*(h_mem32ptr + indx) == (ARR_SIZE - indx));
}
// Test writting device pointer
hipHostGetDevicePointer((void**)&device_ptr64, host_ptr64, 0);
hipHostGetDevicePointer((void**)&device_ptr32, host_ptr32, 0);
hipHostGetDevicePointer((void**)&d_mem64ptr, h_mem64ptr, 0);
hipHostGetDevicePointer((void**)&d_mem32ptr, h_mem32ptr, 0);
// Reset values
*host_ptr64 = 0x0;
@@ -235,10 +254,19 @@ void testWrite() {
HIPCHECK(hipStreamWriteValue64(stream, device_ptr64, value64, writeFlag));
HIPCHECK(hipStreamWriteValue32(stream, device_ptr32, value32, writeFlag));
for (int indx = 0; indx < ARR_SIZE; indx++) {
HIPCHECK(hipStreamWriteValue64(stream, d_mem64ptr + indx, indx, writeFlag));
HIPCHECK(hipStreamWriteValue32(stream, d_mem32ptr + indx, indx, writeFlag));
}
hipStreamSynchronize(stream);
HIPASSERT(*host_ptr64 == value64);
HIPASSERT(*host_ptr32 == value32);
for (int indx = 0; indx < ARR_SIZE ; indx++) {
HIPASSERT(*(h_mem64ptr + indx) == indx);
HIPASSERT(*(h_mem32ptr + indx) == indx);
}
// Test Writing to Signal Memory
HIPCHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag));