AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
Change-Id: Ib6b22852155bd6d448ddd07ac406a0f36732c5fb


[ROCm/hip commit: 80a14899a1]
This commit is contained in:
Jenkins
2022-10-18 07:10:18 -04:00
parent 797a4342d3 3fd7a5dc21
commit 74c170145f
32 changed files with 2223 additions and 122 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/bin/hipcc
+1 -1
View File
@@ -35,7 +35,7 @@ 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;
$arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\ ]/\\$&/g;
}
}
projects/hip/bin/hipcc.pl
-6
View File
@@ -184,9 +184,6 @@ if ($HIP_PLATFORM eq "amd") {
$HIP_CLANG_TARGET = `$HIPCC -print-target-triple`;
chomp($HIP_CLANG_TARGET);
if (! defined $HIP_CLANG_INCLUDE_PATH) {
$HIP_CLANG_INCLUDE_PATH = abs_path("$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include");
}
if (! defined $HIP_INCLUDE_PATH) {
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
}
@@ -199,15 +196,12 @@ if ($HIP_PLATFORM eq "amd") {
print ("HIP_ROCCLR_HOME=$HIP_ROCCLR_HOME\n");
}
print ("HIP_CLANG_PATH=$HIP_CLANG_PATH\n");
print ("HIP_CLANG_INCLUDE_PATH=$HIP_CLANG_INCLUDE_PATH\n");
print ("HIP_INCLUDE_PATH=$HIP_INCLUDE_PATH\n");
print ("HIP_LIB_PATH=$HIP_LIB_PATH\n");
print ("DEVICE_LIB_PATH=$DEVICE_LIB_PATH\n");
print ("HIP_CLANG_TARGET=$HIP_CLANG_TARGET\n");
}
$HIPCXXFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\"";
$HIPCFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\"";
$HIPLDFLAGS .= " -L\"$HIP_LIB_PATH\"";
if ($isWindows) {
$HIPLDFLAGS .= " -lamdhip64";
projects/hip/docs/markdown/hip_porting_guide.md
+15 -12
View File
@@ -468,31 +468,34 @@ int main()
## CU_POINTER_ATTRIBUTE_MEMORY_TYPE
To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'memoryType' as member variable. 'memoryType' indicates input pointer is allocated on device or host.
To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'type' as member variable. 'type' indicates input pointer is allocated on device or host.
For example:
```
double * ptr;
hipMalloc(reinterpret_cast<void**>(&ptr), sizeof(double));
hipPointerAttribute_t attr;
hipPointerGetAttributes(&attr, ptr); /*attr.memoryType will have value as hipMemoryTypeDevice*/
hipPointerGetAttributes(&attr, ptr); /*attr.type will have value as hipMemoryTypeDevice*/
double* ptrHost;
hipHostMalloc(&ptrHost, sizeof(double));
hipPointerAttribute_t attr;
hipPointerGetAttributes(&attr, ptrHost); /*attr.memoryType will have value as hipMemoryTypeHost*/
hipPointerGetAttributes(&attr, ptrHost); /*attr.type will have value as hipMemoryTypeHost*/
```
Please note, hipMemoryType enum values are different from cudaMemoryType enum values.
For example, on AMD platform, memoryType is defined in hip_runtime_api.h,
For example, on AMD platform, hipMemoryType is defined in hip_runtime_api.h,
```
typedef enum hipMemoryType {
hipMemoryTypeHost, ///< Memory is physically located on host
hipMemoryTypeDevice, ///< Memory is physically located on device.
hipMemoryTypeArray, ///< Array memory, physically located on device.
hipMemoryTypeUnified ///< Not used currently
hipMemoryTypeHost = 0, ///< Memory is physically located on host
hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for specific device)
hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for specific device)
hipMemoryTypeUnified = 3, ///< Not used currently
hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified memory system
} hipMemoryType;
Looking into CUDA toolkit, it defines memoryType as following,
```
Looking into CUDA toolkit, it defines cudaMemoryType as following,
```
enum cudaMemoryType
{
cudaMemoryTypeUnregistered = 0, // Unregistered memory.
@@ -500,8 +503,8 @@ enum cudaMemoryType
cudaMemoryTypeDevice = 2, // Device memory.
cudaMemoryTypeManaged = 3, // Managed memory
}
In this case, memoryType translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h.
```
In this case, memory type translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h.
So in any HIP applications which use HIP APIs involving memory types, developers should use #ifdef in order to assign the correct enum values depending on Nvidia or AMD platform.
projects/hip/include/hip/hip_runtime_api.h
+160 -11
View File
@@ -153,24 +153,31 @@ typedef struct hipDeviceProp_t {
} hipDeviceProp_t;
/**
* Memory type (for pointer attributes)
/*
* @brief HIP Memory type (for pointer attributes)
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemoryType {
hipMemoryTypeHost, ///< Memory is physically located on host
hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific
///< device)
hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific
///< device)
hipMemoryTypeUnified, ///< Not used currently
hipMemoryTypeManaged ///< Managed memory, automaticallly managed by the unified memory system
hipMemoryTypeHost = 0, ///< Memory is physically located on host
hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for
///< specific device)
hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for
///< specific device)
hipMemoryTypeUnified = 3, ///< Not used currently
hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified
///< memory system
} hipMemoryType;
/**
* Pointer attributes
*/
typedef struct hipPointerAttribute_t {
enum hipMemoryType memoryType;
union {
// Deprecated, use instead type
enum hipMemoryType memoryType;
enum hipMemoryType type;
};
int device;
void* devicePointer;
void* hostPointer;
@@ -6385,7 +6392,7 @@ hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, un
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count);
hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count __dparm(1));
// doxygen end graph API
/**
* @}
@@ -6768,6 +6775,148 @@ inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
numBlocks, reinterpret_cast<const void*>(f), blockSize, dynSharedMemPerBlk, flags);
}
/**
* @brief Returns grid and block size that achieves maximum potential occupancy for a device function
*
* Returns in \p *min_grid_size and \p *block_size a suggested grid /
* block size pair that achieves the best potential occupancy
* (i.e. the maximum number of active warps on the current device with the smallest number
* of blocks for a particular function).
*
* @param [out] min_grid_size minimum grid size needed to achieve the best potential occupancy
* @param [out] block_size block size required for the best potential occupancy
* @param [in] func device function symbol
* @param [in] block_size_to_dynamic_smem_size - a unary function/functor that takes block size,
* and returns the size, in bytes, of dynamic shared memory needed for a block
* @param [in] block_size_limit the maximum block size \p func is designed to work with. 0 means no limit.
* @param [in] flags reserved
*
* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue,
* #hipErrorUnknown
*/
template<typename UnaryFunction, class T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(
int* min_grid_size,
int* block_size,
T func,
UnaryFunction block_size_to_dynamic_smem_size,
int block_size_limit = 0,
unsigned int flags = 0) {
if (min_grid_size == nullptr || block_size == nullptr ||
reinterpret_cast<const void*>(func) == nullptr) {
return hipErrorInvalidValue;
}
int dev;
hipError_t status;
if ((status = hipGetDevice(&dev)) != hipSuccess) {
return status;
}
int max_threads_per_cu;
if ((status = hipDeviceGetAttribute(&max_threads_per_cu,
hipDeviceAttributeMaxThreadsPerMultiProcessor, dev)) != hipSuccess) {
return status;
}
int warp_size;
if ((status = hipDeviceGetAttribute(&warp_size,
hipDeviceAttributeWarpSize, dev)) != hipSuccess) {
return status;
}
int max_cu_count;
if ((status = hipDeviceGetAttribute(&max_cu_count,
hipDeviceAttributeMultiprocessorCount, dev)) != hipSuccess) {
return status;
}
struct hipFuncAttributes attr;
if ((status = hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(func))) != hipSuccess) {
return status;
}
// Initial limits for the execution
const int func_max_threads_per_block = attr.maxThreadsPerBlock;
if (block_size_limit == 0) {
block_size_limit = func_max_threads_per_block;
}
if (func_max_threads_per_block < block_size_limit) {
block_size_limit = func_max_threads_per_block;
}
const int block_size_limit_aligned =
((block_size_limit + (warp_size - 1)) / warp_size) * warp_size;
// For maximum search
int max_threads = 0;
int max_block_size{};
int max_num_blocks{};
for (int block_size_check_aligned = block_size_limit_aligned;
block_size_check_aligned > 0;
block_size_check_aligned -= warp_size) {
// Make sure the logic uses the requested limit and not aligned
int block_size_check = (block_size_limit < block_size_check_aligned) ?
block_size_limit : block_size_check_aligned;
size_t dyn_smem_size = block_size_to_dynamic_smem_size(block_size_check);
int optimal_blocks;
if ((status = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
&optimal_blocks, func, block_size_check, dyn_smem_size, flags)) != hipSuccess) {
return status;
}
int total_threads = block_size_check * optimal_blocks;
if (total_threads > max_threads) {
max_block_size = block_size_check;
max_num_blocks = optimal_blocks;
max_threads = total_threads;
}
// Break if the logic reached possible maximum
if (max_threads_per_cu == max_threads) {
break;
}
}
// Grid size is the number of blocks per CU * CU count
*min_grid_size = max_num_blocks * max_cu_count;
*block_size = max_block_size;
return status;
}
/**
* @brief Returns grid and block size that achieves maximum potential occupancy for a device function
*
* Returns in \p *min_grid_size and \p *block_size a suggested grid /
* block size pair that achieves the best potential occupancy
* (i.e. the maximum number of active warps on the current device with the smallest number
* of blocks for a particular function).
*
* @param [out] min_grid_size minimum grid size needed to achieve the best potential occupancy
* @param [out] block_size block size required for the best potential occupancy
* @param [in] func device function symbol
* @param [in] block_size_to_dynamic_smem_size - a unary function/functor that takes block size,
* and returns the size, in bytes, of dynamic shared memory needed for a block
* @param [in] block_size_limit the maximum block size \p func is designed to work with. 0 means no limit.
*
* @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue,
* #hipErrorUnknown
*/
template<typename UnaryFunction, class T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMem(
int* min_grid_size,
int* block_size,
T func,
UnaryFunction block_size_to_dynamic_smem_size,
int block_size_limit = 0)
{
return hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(min_grid_size, block_size, func,
block_size_to_dynamic_smem_size, block_size_limit);
}
template <typename F>
inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json
+19 -15
View File
@@ -6,35 +6,39 @@
"Unit_hipGraphAddHostNode_ClonedGraphwithHostNode",
"Unit_hipEventIpc",
"Unit_hipMalloc3D_Negative",
"Unit_hipStreamValue_Write",
"Unit_hipMemPoolApi_BasicAlloc",
"Unit_hipMemPoolApi_BasicTrim",
"Unit_hipMemPoolApi_BasicReuse",
"Unit_hipMemPoolApi_Opportunistic",
"Unit_hipGraphMemcpyNodeSetParams_Functional",
"Unit_hipMalloc3D_ValidatePitch",
"Unit_hipArrayCreate_happy",
"Unit_hipMemAllocPitch_ValidatePitch",
"Unit_hipArrayCreate_happy - int",
"Unit_hipArrayCreate_happy - int4",
"Unit_hipArrayCreate_happy - short2",
"Unit_hipArrayCreate_happy - char",
"Unit_hipArrayCreate_happy - char4",
"Unit_hipArrayCreate_happy - float",
"Unit_hipArrayCreate_happy - float2",
"Unit_hipArrayCreate_happy - float4",
"Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional",
"Unit_hipMallocManaged_MultiChunkMultiDevice",
"Unit_hipMallocManaged_TwoPointers - int",
"Unit_hipMallocManaged_TwoPointers - float",
"Unit_hipMallocManaged_TwoPointers - double",
"Unit_hipMallocManaged_OverSubscription",
"Unit_hipMallocManaged_CoherentTstWthAdvise",
"Unit_hipMallocManaged_Advanced",
"Unit_hipMemRangeGetAttribute_TstCountParam",
"Unit_hipMemRangeGetAttribute_NegativeTests",
"Unit_hipMemRangeGetAttribute_AccessedBy1",
"Unit_hipMemRangeGetAttribte_3",
"Unit_hipMemRangeGetAttribute_4",
"Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr",
"Unit_hipMemAdvise_TstFlags",
"Unit_hipMemAdvise_PrefrdLoc",
"Unit_hipMemAdvise_ReadMostly",
"Unit_hipMemAdvise_TstFlgOverrideEffect",
"Unit_hipMemAdvise_TstAccessedByFlg",
"Unit_hipMemAdvise_TstAccessedByFlg4",
"Unit_hipMemAdvise_TstMemAdvisePrefrdLoc",
"Unit_hipMemAdvise_TstMemAdviseMultiFlag",
"Unit_hipMemAdvise_ReadMosltyMgpuTst",
"Unit_hipMemAdvise_TstSetUnsetPrfrdLoc",
"Unit_hipMallocManaged_DeviceContextChange - unsigned char",
"Unit_hipMallocManaged_DeviceContextChange - int",
"Unit_hipMallocManaged_DeviceContextChange - float",
"Unit_hipMallocManaged_DeviceContextChange - double",
"Unit_hipGraphNodeGetDependentNodes_Functional",
"Unit_hipGraphNodeGetDependencies_Functional",
"Unit_hipMemGetInfo_Negative",
"Unit_hipStreamCreateWithPriority_ValidateWithEvents",
"Unit_hipStreamPerThread_StrmWaitEvt",
"Unit_hipMemGetInfo_DifferentMallocSmall",
projects/hip/tests/catch/include/resource_guards.hh
+125
View File
@@ -0,0 +1,125 @@
/*
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 <hip/hip_runtime_api.h>
enum class LinearAllocs {
malloc,
mallocAndRegister,
hipHostMalloc,
hipMalloc,
hipMallocManaged,
};
template <typename T> class LinearAllocGuard {
public:
LinearAllocGuard(const LinearAllocs allocation_type, const size_t size,
const unsigned int flags = 0u)
: allocation_type_{allocation_type} {
switch (allocation_type_) {
case LinearAllocs::malloc:
ptr_ = host_ptr_ = reinterpret_cast<T*>(malloc(size));
break;
case LinearAllocs::mallocAndRegister:
host_ptr_ = reinterpret_cast<T*>(malloc(size));
HIP_CHECK(hipHostRegister(host_ptr_, size, flags));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&ptr_), host_ptr_, 0u));
break;
case LinearAllocs::hipHostMalloc:
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr_), size, flags));
host_ptr_ = ptr_;
break;
case LinearAllocs::hipMalloc:
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr_), size));
break;
case LinearAllocs::hipMallocManaged:
HIP_CHECK(hipMallocManaged(reinterpret_cast<void**>(&ptr_), size, flags ? flags : 1u));
host_ptr_ = ptr_;
}
}
LinearAllocGuard(const LinearAllocGuard&) = delete;
LinearAllocGuard(LinearAllocGuard&&) = delete;
~LinearAllocGuard() {
// No Catch macros, don't want to possibly throw in the destructor
switch (allocation_type_) {
case LinearAllocs::malloc:
free(ptr_);
break;
case LinearAllocs::mallocAndRegister:
// Cast to void to suppress nodiscard warnings
static_cast<void>(hipHostUnregister(host_ptr_));
free(host_ptr_);
break;
case LinearAllocs::hipHostMalloc:
static_cast<void>(hipHostFree(ptr_));
break;
case LinearAllocs::hipMalloc:
case LinearAllocs::hipMallocManaged:
static_cast<void>(hipFree(ptr_));
}
}
T* ptr() { return ptr_; };
T* const ptr() const { return ptr_; };
T* host_ptr() { return host_ptr_; }
T* const host_ptr() const { return host_ptr(); }
private:
const LinearAllocs allocation_type_;
T* ptr_ = nullptr;
T* host_ptr_ = nullptr;
};
enum class Streams { nullstream, perThread, created };
class StreamGuard {
public:
StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
switch (stream_type_) {
case Streams::nullstream:
stream_ = nullptr;
break;
case Streams::perThread:
stream_ = hipStreamPerThread;
break;
case Streams::created:
HIP_CHECK(hipStreamCreate(&stream_));
}
}
StreamGuard(const StreamGuard&) = delete;
StreamGuard(StreamGuard&&) = delete;
~StreamGuard() {
if (stream_type_ == Streams::created) {
static_cast<void>(hipStreamDestroy(stream_));
}
}
hipStream_t stream() const { return stream_; }
private:
const Streams stream_type_;
hipStream_t stream_;
};
projects/hip/tests/catch/include/utils.hh
+102
View File
@@ -0,0 +1,102 @@
/*
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 <chrono>
#include <hip_test_common.hh>
#include <hip/hip_runtime_api.h>
namespace {
inline constexpr size_t kPageSize = 4096;
} // anonymous namespace
template <typename T>
void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) {
const auto ret = std::mismatch(expected, expected + num_elements, actual);
if (ret.first != expected + num_elements) {
const auto idx = std::distance(expected, ret.first);
INFO("Value mismatch at index: " << idx);
REQUIRE(expected[idx] == actual[idx]);
}
}
template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
const auto it = std::find_if_not(
begin, end, [expected_value](const int elem) { return expected_value == elem; });
if (it != end) {
const auto idx = std::distance(begin, it);
INFO("Value mismatch at index " << idx);
REQUIRE(expected_value == *it);
}
}
template <typename T>
void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) {
ArrayFindIfNot(array, array + num_elements, expected_value);
}
template <typename T>
__global__ void VectorIncrement(T* const vec, const T increment_value, 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) {
vec[i] += increment_value;
}
}
template <typename T> __global__ void VectorSet(T* const vec, const T value, 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) {
vec[i] = value;
}
}
// Will execute for atleast interval milliseconds
static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) {
while (interval--) {
uint64_t start = clock();
while (clock() - start < ticks_per_ms) {
}
}
}
inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) {
int ticks_per_ms = 0;
// Clock rate is in kHz => number of clock ticks in a millisecond
HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0));
Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms);
HIP_CHECK(hipGetLastError());
}
template <typename... Attributes>
inline bool DeviceAttributesSupport(const int device, Attributes... attributes) {
constexpr auto DeviceAttributeSupport = [](const int device,
const hipDeviceAttribute_t attribute) {
int value = 0;
HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device));
return value;
};
return (... && DeviceAttributeSupport(device, attributes));
}
projects/hip/tests/catch/multiproc/CMakeLists.txt
+12
View File
@@ -15,10 +15,22 @@ set(LINUX_TEST_SRC
hipIpcEventHandle.cc
hipIpcMemAccessTest.cc
deviceAllocationMproc.cc
hipNoGpuTsts.cc
)
add_custom_target(dummy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/dummy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../multiproc/dummy_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
# the last argument linker libraries is required for this test but optional to the function
if(HIP_PLATFORM MATCHES "nvidia")
hip_add_exe_to_target(NAME MultiProc
TEST_SRC ${LINUX_TEST_SRC}
TEST_TARGET_NAME build_tests
LINKER_LIBS nvrtc)
elseif(HIP_PLATFORM MATCHES "amd")
hip_add_exe_to_target(NAME MultiProc
TEST_SRC ${LINUX_TEST_SRC}
TEST_TARGET_NAME build_tests
LINKER_LIBS ${CMAKE_DL_LIBS})
endif()
add_dependencies(build_tests dummy_kernel.code)
projects/hip/tests/catch/multiproc/dummy_kernel.cpp
+26
View File
@@ -0,0 +1,26 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
extern "C" __global__ void dummy_ker() {
}
projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc
+1687
View File
File diff suppressed because it is too large Load Diff
projects/hip/tests/catch/unit/memory/CMakeLists.txt
+1 -2
View File
@@ -183,5 +183,4 @@ endif()
hip_add_exe_to_target(NAME MemoryTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests
COMPILE_OPTIONS -std=c++14)
TEST_TARGET_NAME build_tests)
projects/hip/tests/catch/unit/memory/MemUtils.hh
+2 -2
View File
@@ -344,10 +344,10 @@ static inline void memsetCheck(T* aPtr, size_t value, memType memType, MultiDDat
template <typename T> static inline void freeStuff(T* aPtr, allocType type) {
switch (type) {
case allocType::deviceMalloc:
hipFree(aPtr);
HIP_CHECK(hipFree(aPtr));
break;
case allocType::hostMalloc:
hipHostFree(aPtr);
HIP_CHECK(hipHostFree(aPtr));
break;
default: // for host and device registered
HIP_CHECK(hipHostUnregister(aPtr));
projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc
+2 -2
View File
@@ -501,8 +501,8 @@ void DrvMemcpy3D<T>::HostArray_DrvMemcpy3D(bool device_context_change) {
/* DeAllocating the memory */
template <typename T>
void DrvMemcpy3D<T>::DeAllocateMemory() {
hipArrayDestroy(arr);
hipArrayDestroy(arr1);
HIP_CHECK(hipArrayDestroy(arr));
HIP_CHECK(hipArrayDestroy(arr1));
free(hData);
}
projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc
+1 -1
View File
@@ -93,7 +93,7 @@ TEST_CASE("Unit_hipMallocManaged_HostDeviceConcurrent") {
std::thread host_thread(HostKernelDouble, Hmm, hPtr, N);
KernelDouble<<<dim3(blocks), dim3(threadsPerBlock), 0, 0>>>(Hmm, dPtr, N);
host_thread.join();
hipMemcpy(resPtr, dPtr, N * sizeof(float), hipMemcpyDeviceToHost);
HIP_CHECK(hipMemcpy(resPtr, dPtr, N * sizeof(float), hipMemcpyDeviceToHost));
for (size_t i = 0; i < N; i++) {
REQUIRE(hPtr[i] == resPtr[i]);
projects/hip/tests/catch/unit/memory/hipMemcpy.cc
+5 -5
View File
@@ -576,7 +576,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch",
HIP_CHECK(hipSetDevice(HipTest::RAND_R(&seed) % (numDevices-1)+1));
int device;
hipGetDevice(&device);
HIP_CHECK(hipGetDevice(&device));
std::cout <<"hipMemcpy is set to happen between device 0 and device "
<<device << std::endl;
HipTest::initArrays<TestType>(&X_d, &Y_d, &Z_d, nullptr,
@@ -588,10 +588,10 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch",
C_h[j] = 0;
}
hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice);
hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice);
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc
+1 -1
View File
@@ -135,7 +135,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_multiDevice-D2D", ""
size_t width{NUM_W * sizeof(TestType)};
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc
+2 -2
View File
@@ -160,7 +160,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-Host&PinnedMem", ""
hipStream_t stream;
if (numDevices > 1) {
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipStreamCreate(&stream));
@@ -250,7 +250,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", ""
hipStream_t stream;
if (numDevices > 1) {
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc
+2 -2
View File
@@ -190,7 +190,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
@@ -246,7 +246,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_multiDeviceContextChange") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc
+2 -2
View File
@@ -217,7 +217,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDevicePinnedHostMem") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
@@ -277,7 +277,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDeviceContextChange") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc
+2 -2
View File
@@ -150,7 +150,7 @@ TEST_CASE("Unit_hipMemcpy_NullCheck") {
&A_h, &B_h, &C_h,
NUM_ELM*sizeof(float));
hipStream_t stream;
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpy(A_d, C_h,
NUM_ELM*sizeof(float),
hipMemcpyHostToDevice));
@@ -240,7 +240,7 @@ TEST_CASE("Unit_hipMemcpy_HalfMemCopy") {
&A_h, &B_h, &C_h,
NUM_ELM*sizeof(float));
hipStream_t stream;
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
SECTION("hipMemcpyHtoD half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc
+6 -6
View File
@@ -112,7 +112,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float,
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipSetDevice(0));
hipStream_t stream;
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false);
@@ -147,7 +147,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int,
TestType *A_Ph{nullptr}, *B_Ph{nullptr};
HIP_CHECK(hipSetDevice(0));
hipStream_t stream;
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
HipTest::initArrays<TestType>(&A_d, &B_d, nullptr,
&A_h, &B_h, nullptr,
NUM_ELM*sizeof(TestType));
@@ -373,10 +373,10 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch",
C_h[j] = 0;
}
hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream);
hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream);
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream));
HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc
+2 -2
View File
@@ -35,7 +35,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Negative") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
// Initialization of variables
int *A_d{nullptr}, *B_d{nullptr};
@@ -102,7 +102,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Basic") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc
+3 -3
View File
@@ -39,7 +39,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Negative") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
// Initialization of variables
int *A_d{nullptr}, *B_d{nullptr};
@@ -116,7 +116,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
// Initialization of Variables on GPU-0
int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
@@ -202,7 +202,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_StreamOnDiffDevice") {
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, 0, 1);
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc
+4 -4
View File
@@ -105,7 +105,7 @@ void memcpyTests<T>::Memcpy_And_verify(bool *ret_val) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDefault));
@@ -155,7 +155,7 @@ void memcpyTests<T>::Memcpy_And_verify(bool *ret_val) {
int canAccessPeer = 0;
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(T)));
@@ -198,7 +198,7 @@ void memcpyTests<T>::Memcpy_And_verify(bool *ret_val) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i],
NUM_ELM * sizeof(T),
@@ -252,7 +252,7 @@ void memcpyTests<T>::Memcpy_And_verify(bool *ret_val) {
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j));
if (canAccessPeer) {
HIPCHECK(hipSetDevice(j));
HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]),
projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc
+3 -3
View File
@@ -44,8 +44,8 @@ __global__ void Kernel(float* hostRes, int clkRate) {
TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
int blocks = 2;
float* hostRes;
hipHostMalloc(&hostRes, blocks * sizeof(float),
hipHostMallocMapped);
HIP_CHECK(hipHostMalloc(&hostRes, blocks * sizeof(float),
hipHostMallocMapped));
hostRes[0] = 0;
hostRes[1] = 0;
int clkRate;
@@ -60,6 +60,6 @@ TEST_CASE("Unit_hipHostMalloc_CoherentAccess") {
while (hostRes[eleCounter] == 0) {printf("waiting for counter inc\n");}
eleCounter++;
}
hipHostFree(reinterpret_cast<void *>(hostRes));
HIP_CHECK(hipHostFree(reinterpret_cast<void *>(hostRes)));
}
projects/hip/tests/catch/unit/memory/hipMemset.cc
+11 -11
View File
@@ -260,21 +260,21 @@ TEST_CASE("Unit_hipMemset_2AsyncOperations") {
std::vector<float> v;
v.resize(2048);
float* p2, *p3;
hipMalloc(reinterpret_cast<void**>(&p2), 4096 + 4096*2);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&p2), 4096 + 4096*2));
p3 = p2+2048;
hipStream_t s;
hipStreamCreate(&s);
hipMemsetAsync(p2, 0, 32*32*4, s);
hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
hipStreamSynchronize(s);
HIP_CHECK(hipStreamCreate(&s));
HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s));
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s));
HIP_CHECK(hipStreamSynchronize(s));
for (int i = 0; i < 256; ++i) {
hipMemsetAsync(p2, 0, 32*32*4, s);
hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s);
HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s));
HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s));
}
hipStreamSynchronize(s);
hipDeviceSynchronize();
hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost);
hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost);
HIP_CHECK(hipStreamSynchronize(s));
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost));
REQUIRE(v[0] == 0);
REQUIRE(v[1024] == 1.75f);
projects/hip/tests/catch/unit/memory/hipMemset2D.cc
+3 -3
View File
@@ -78,7 +78,7 @@ TEST_CASE("Unit_hipMemset2D_BasicFunctional") {
}
}
hipFree(A_d);
HIP_CHECK(hipFree(A_d));
free(A_h);
}
@@ -120,7 +120,7 @@ TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") {
}
}
hipFree(A_d);
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipStreamDestroy(stream));
free(A_h);
}
@@ -169,7 +169,7 @@ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") {
}
}
hipFree(A_d);
HIP_CHECK(hipFree(A_d));
free(A_h);
}
projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc
+5 -5
View File
@@ -83,7 +83,7 @@ void checkMemset(T value, size_t count, MemsetType memsetType, bool async = fals
MemsetMallocType mallocType = hipDeviceMalloc_t, T* devPtr = nullptr) {
hipStream_t stream{nullptr};
if (async) {
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
}
// Allocate Memory
@@ -254,7 +254,7 @@ template <typename T>
void checkMemset2D(T value, size_t width, size_t height, bool async = false, size_t pitch = 0,
T* devPtr = nullptr) {
hipStream_t stream{nullptr};
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
constexpr size_t elementSize = sizeof(T);
bool freeDevPtr = false;
if (devPtr == nullptr) {
@@ -277,7 +277,7 @@ void checkMemset2D(T value, size_t width, size_t height, bool async = false, siz
if (freeDevPtr) {
HIP_CHECK(hipFree(devPtr));
}
hipStreamDestroy(stream);
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_2D") {
@@ -430,7 +430,7 @@ void check_device_data_3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent exten
template <typename T>
void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool async = false) {
hipStream_t stream{nullptr};
hipStreamCreate(&stream);
HIP_CHECK(hipStreamCreate(&stream));
if (devPitchedPtr.ptr == nullptr) {
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
}
@@ -445,7 +445,7 @@ void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool
if (extent.width * extent.height * extent.depth > 0) {
check_device_data_3D(devPitchedPtr, value, extent);
}
hipStreamDestroy(stream);
HIP_CHECK(hipStreamDestroy(stream));
}
void check_memset_3D(std::string sectionStr, size_t width, size_t height, size_t depth,
projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc
+1 -1
View File
@@ -105,7 +105,7 @@ TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") {
size_t pitch_A;
HIP_CHECK(hipMallocPitch(&A_d, &pitch_A, width, height));
testHipMemset2DApis(dst, pitch_A, memsetVal, width, height);
hipFree(A_d);
HIP_CHECK(hipFree(A_d));
}
TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") {
projects/hip/tests/catch/unit/memory/hipMemsetSync.cc
+2 -2
View File
@@ -379,10 +379,10 @@ void memsetCheck(T* aPtr, size_t value, memSetType memsetType, MultiDData& data,
template <typename T> void freeStuff(T* aPtr, allocType type) {
switch (type) {
case allocType::deviceMalloc:
hipFree(aPtr);
HIP_CHECK(hipFree(aPtr));
break;
case allocType::hostMalloc:
hipHostFree(aPtr);
HIP_CHECK(hipHostFree(aPtr));
break;
case allocType::hostRegisted:
HIP_CHECK(hipHostUnregister(aPtr));
projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc
+9 -9
View File
@@ -43,7 +43,7 @@ constexpr size_t N{1000000};
bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) {
return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) &&
(lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) &&
(lhs.type == rhs.type) && (lhs.device == rhs.device) &&
(lhs.allocationFlags == rhs.allocationFlags));
}
@@ -68,7 +68,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) {
void resetAttribs(hipPointerAttribute_t* attribs) {
attribs->hostPointer = reinterpret_cast<void*>(-1);
attribs->devicePointer = reinterpret_cast<void*>(-1);
attribs->memoryType = hipMemoryTypeHost;
attribs->type = hipMemoryTypeHost;
attribs->device = -2;
attribs->isManaged = -1;
attribs->allocationFlags = 0xffff;
@@ -77,9 +77,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) {
void printAttribs(const hipPointerAttribute_t* attribs) {
printf(
"hostPointer:%p devicePointer:%p memType:%s deviceId:%d isManaged:%d "
"hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d "
"allocationFlags:%u\n",
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType),
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type),
attribs->device, attribs->isManaged, attribs->allocationFlags);
}
@@ -151,14 +151,14 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
if (isDevice) {
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptr), reference[i]._sizeBytes));
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
reference[i]._attrib.type = hipMemoryTypeDevice;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = NULL;
reference[i]._attrib.allocationFlags = 0;
} else {
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&ptr), reference[i]._sizeBytes,
hipHostMallocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeHost;
reference[i]._attrib.type = hipMemoryTypeHost;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = ptr;
reference[i]._attrib.allocationFlags = 0;
@@ -190,7 +190,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
checkPointer(ref, i, 2, (ptr + ref._sizeBytes - 1));
}
if (ref._attrib.memoryType == hipMemoryTypeDevice) {
if (ref._attrib.type == hipMemoryTypeDevice) {
HIP_CHECK(hipFree(ref._pointer));
} else {
HIP_CHECK(hipHostFree(ref._pointer));
@@ -369,11 +369,11 @@ TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") {
// Memory address and type check
if (MemoryType == MemoryTypes::DeviceMemory) {
REQUIRE(attributes.memoryType == hipMemoryTypeDevice);
REQUIRE(attributes.type == hipMemoryTypeDevice);
REQUIRE(attributes.devicePointer == ptr);
REQUIRE(attributes.hostPointer == nullptr);
} else if (MemoryType == MemoryTypes::HostMemory) {
REQUIRE(attributes.memoryType == hipMemoryTypeHost);
REQUIRE(attributes.type == hipMemoryTypeHost);
REQUIRE(attributes.hostPointer == ptr);
} else if (MemoryType == MemoryTypes::MappedMemory) {
int* devicePtr{nullptr};
projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp
+7 -7
View File
@@ -40,7 +40,7 @@ size_t Nbytes = 0;
bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) {
return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) &&
(lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) &&
(lhs.type == rhs.type) && (lhs.device == rhs.device) &&
(lhs.allocationFlags == rhs.allocationFlags));
};
@@ -65,7 +65,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) {
void resetAttribs(hipPointerAttribute_t* attribs) {
attribs->hostPointer = (void*)(-1);
attribs->devicePointer = (void*)(-1);
attribs->memoryType = hipMemoryTypeHost;
attribs->type = hipMemoryTypeHost;
attribs->device = -2;
attribs->isManaged = -1;
attribs->allocationFlags = 0xffff;
@@ -74,9 +74,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) {
void printAttribs(const hipPointerAttribute_t* attribs) {
printf(
"hostPointer:%p devicePointer:%p memoryType:%s deviceId:%d isManaged:%d "
"hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d "
"allocationFlags:%u\n",
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType),
attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type),
attribs->device, attribs->isManaged, attribs->allocationFlags);
};
@@ -229,13 +229,13 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
if (isDevice) {
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
HIPCHECK(hipMalloc((void**)&ptr, reference[i]._sizeBytes));
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
reference[i]._attrib.type = hipMemoryTypeDevice;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = NULL;
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
} else {
HIPCHECK(hipHostMalloc((void**)&ptr, reference[i]._sizeBytes, hipHostMallocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeHost;
reference[i]._attrib.type = hipMemoryTypeHost;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = ptr;
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
@@ -265,7 +265,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
checkPointer(ref, i, 2, (char*)ref._pointer + ref._sizeBytes - 1);
}
if (ref._attrib.memoryType == hipMemoryTypeDevice) {
if (ref._attrib.type == hipMemoryTypeDevice) {
hipFree(ref._pointer);
} else {
hipHostFree(ref._pointer);