diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 42fa1e8a20..dfb5a8f6b8 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -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; } } diff --git a/projects/hip/bin/hipcc.pl b/projects/hip/bin/hipcc.pl index da9559be9e..606727f4cd 100755 --- a/projects/hip/bin/hipcc.pl +++ b/projects/hip/bin/hipcc.pl @@ -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"; diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 179ad64a2f..fe1972e3b3 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -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(&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. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 1d8a6cc731..aa36787636 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -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(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 +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(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(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 +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 inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json index 48630e0ce1..002666b345 100644 --- a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -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", diff --git a/projects/hip/tests/catch/include/resource_guards.hh b/projects/hip/tests/catch/include/resource_guards.hh new file mode 100644 index 0000000000..7e6179c81a --- /dev/null +++ b/projects/hip/tests/catch/include/resource_guards.hh @@ -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 +#include + +enum class LinearAllocs { + malloc, + mallocAndRegister, + hipHostMalloc, + hipMalloc, + hipMallocManaged, +}; + +template 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(malloc(size)); + break; + case LinearAllocs::mallocAndRegister: + host_ptr_ = reinterpret_cast(malloc(size)); + HIP_CHECK(hipHostRegister(host_ptr_, size, flags)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&ptr_), host_ptr_, 0u)); + break; + case LinearAllocs::hipHostMalloc: + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr_), size, flags)); + host_ptr_ = ptr_; + break; + case LinearAllocs::hipMalloc: + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr_), size)); + break; + case LinearAllocs::hipMallocManaged: + HIP_CHECK(hipMallocManaged(reinterpret_cast(&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(hipHostUnregister(host_ptr_)); + free(host_ptr_); + break; + case LinearAllocs::hipHostMalloc: + static_cast(hipHostFree(ptr_)); + break; + case LinearAllocs::hipMalloc: + case LinearAllocs::hipMallocManaged: + static_cast(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(hipStreamDestroy(stream_)); + } + } + + hipStream_t stream() const { return stream_; } + + private: + const Streams stream_type_; + hipStream_t stream_; +}; \ No newline at end of file diff --git a/projects/hip/tests/catch/include/utils.hh b/projects/hip/tests/catch/include/utils.hh new file mode 100644 index 0000000000..9edffc6f7c --- /dev/null +++ b/projects/hip/tests/catch/include/utils.hh @@ -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 + +#include +#include + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +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 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 +void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) { + ArrayFindIfNot(array, array + num_elements, expected_value); +} + +template +__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 __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 +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)); +} \ No newline at end of file diff --git a/projects/hip/tests/catch/multiproc/CMakeLists.txt b/projects/hip/tests/catch/multiproc/CMakeLists.txt index 128384ea3a..5541067207 100644 --- a/projects/hip/tests/catch/multiproc/CMakeLists.txt +++ b/projects/hip/tests/catch/multiproc/CMakeLists.txt @@ -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) + diff --git a/projects/hip/tests/catch/multiproc/dummy_kernel.cpp b/projects/hip/tests/catch/multiproc/dummy_kernel.cpp new file mode 100644 index 0000000000..876fe4e83e --- /dev/null +++ b/projects/hip/tests/catch/multiproc/dummy_kernel.cpp @@ -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() { +} diff --git a/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc b/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc new file mode 100644 index 0000000000..ad4b9813f4 --- /dev/null +++ b/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc @@ -0,0 +1,1687 @@ +/* +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. +*/ + +#include +#include +#include +#include + +#ifdef __linux__ +#include +#include +#include + +// Common Definitions and Macros +#if HT_AMD + static const char* visibility_env_var = "HIP_VISIBLE_DEVICES"; +#else + static const char* visibility_env_var = "CUDA_VISIBLE_DEVICES"; +#endif + +#define UNUSED(expr) do { (void)(expr); } while (0) + +void Callback(hipStream_t stream, hipError_t status, + void* userData) { + UNUSED(stream); + HIP_CHECK(status); + REQUIRE(userData == NULL); +} + +// Dummy RTC function +static constexpr auto rtcfn { + R"( +extern "C" +__global__ void fn() +{ +} +)"}; + +// Dummy Kernel Function +static __global__ void fn() { +} + +// No GPU test case for hipGetDeviceCount +static bool NoGpuTst_hipGetDeviceCount() { + bool passed = false; + hipError_t err; + int nGpus = 0; + err = hipGetDeviceCount(&nGpus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGetDeviceCount: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipSetDevice +static bool NoGpuTst_hipSetDevice() { + bool passed = false; + hipError_t err; + err = hipSetDevice(0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipSetDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetAttribute +static bool NoGpuTst_hipDeviceGetAttribute() { + bool passed = false; + hipError_t err; + int pi = 0; + hipDeviceAttribute_t attr = hipDeviceAttributeCudaCompatibleBegin; + err = hipDeviceGetAttribute(&pi, attr, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetAttribute: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetP2PAttribute +static bool NoGpuTst_hipDeviceGetP2PAttribute() { + bool passed = false; + hipError_t err; + int srcDevice = 0, dstDevice = 1; + int value{-1}; + err = hipDeviceGetP2PAttribute(&value, hipDevP2PAttrPerformanceRank, + srcDevice, dstDevice); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetP2PAttribute: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetPCIBusId +static bool NoGpuTst_hipDeviceGetPCIBusId() { + bool passed = false; + hipError_t err; + int device = 0; + char pciBusId[128]; + err = hipDeviceGetPCIBusId(pciBusId, 128, device); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetPCIBusId: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetByPCIBusId +static bool NoGpuTst_hipDeviceGetByPCIBusId() { + bool passed = false; + hipError_t err; + int device = 0; + char pciBusId[128]; + err = hipDeviceGetByPCIBusId(&device, pciBusId); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetByPCIBusId: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceTotalMem +static bool NoGpuTst_hipDeviceTotalMem() { + bool passed = false; + hipError_t err, err_exp; +#if HT_AMD + err_exp = hipErrorNoDevice; +#else + err_exp = hipErrorNotInitialized; +#endif + int device = 0; + size_t bytes = 0; + err = hipDeviceTotalMem(&bytes, device); + if (err == err_exp) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceTotalMem: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipFuncSetSharedMemConfig +static bool NoGpuTst_hipFuncSetSharedMemConfig() { + bool passed = false; + hipError_t err; + err = hipFuncSetSharedMemConfig(reinterpret_cast(&fn), + hipSharedMemBankSizeDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamCreate +static bool NoGpuTst_hipStreamCreate() { + bool passed = false; + hipError_t err; + hipStream_t stream = nullptr; + err = hipStreamCreate(&stream); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamCreate: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamCreateWithPriority +static bool NoGpuTst_hipStreamCreateWithPriority() { + bool passed = false; + hipError_t err; + hipStream_t mystream = nullptr; + err = hipStreamCreateWithPriority(&mystream, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamCreateWithPriority: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamSynchronize +static bool NoGpuTst_hipStreamSynchronize() { + bool passed = false; + hipError_t err; + err = hipStreamSynchronize(0); // sync with null stream + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamSynchronize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamGetFlags +static bool NoGpuTst_hipStreamGetFlags() { + bool passed = false; + hipError_t err; + unsigned int flags = 0; + err = hipStreamGetFlags(0, &flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamGetFlags: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test case for hipExtStreamCreateWithCUMask +static bool NoGpuTst_hipExtStreamCreateWithCUMask() { + bool passed = false; + hipError_t err; + hipStream_t mystream = nullptr; + uint32_t cuMaskSize = 0; + const uint32_t cuMask = 0; + err = hipExtStreamCreateWithCUMask(&mystream, cuMaskSize, &cuMask); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtStreamCreateWithCUMask: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipStreamAddCallback +static bool NoGpuTst_hipStreamAddCallback() { + bool passed = false; + hipError_t err; + void* userData = NULL; + unsigned int flags = 0; + err = hipStreamAddCallback(0, Callback, userData, flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamAddCallback: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipEventCreateWithFlags +static bool NoGpuTst_hipEventCreateWithFlags() { + bool passed = false; + hipError_t err; + hipEvent_t event = 0; + err = hipEventCreateWithFlags(&event, hipEventDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipEventCreateWithFlags: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceSynchronize +static bool NoGpuTst_hipDeviceSynchronize() { + bool passed = false; + hipError_t err; + err = hipDeviceSynchronize(); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceSynchronize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipGetDevice +static bool NoGpuTst_hipGetDevice() { + bool passed = false; + hipError_t err; + int deviceId = 0; + err = hipGetDevice(&deviceId); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGetDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetCacheConfig +static bool NoGpuTst_hipDeviceGetCacheConfig() { + bool passed = false; + hipError_t err; + hipFuncCache_t cacheConfig = hipFuncCachePreferNone; + err = hipDeviceGetCacheConfig(&cacheConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetCacheConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetSharedMemConfig +static bool NoGpuTst_hipDeviceGetSharedMemConfig() { + bool passed = false; + hipError_t err; + hipSharedMemConfig pConfig = hipSharedMemBankSizeDefault; + err = hipDeviceGetSharedMemConfig(&pConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceSetSharedMemConfig +static bool NoGpuTst_hipDeviceSetSharedMemConfig() { + bool passed = false; + hipError_t err; + hipSharedMemConfig pConfig = hipSharedMemBankSizeDefault; + err = hipDeviceSetSharedMemConfig(pConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceSetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipPointerGetAttributes +static bool NoGpuTst_hipPointerGetAttributes() { + bool passed = false; + hipError_t err; + hipPointerAttribute_t attributes = {}; + char* A_d; + err = hipPointerGetAttributes(&attributes, A_d); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipPointerGetAttributes: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test case for hipExtMallocWithFlags +static bool NoGpuTst_hipExtMallocWithFlags() { + bool passed = false; + hipError_t err; + int *Ptr = nullptr; + unsigned int flags = 0; + err = hipExtMallocWithFlags(reinterpret_cast(Ptr), 128, flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtMallocWithFlags: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipMallocManaged +static bool NoGpuTst_hipMallocManaged() { + bool passed = false; + hipError_t err; + float *dev_ptr; + err = hipMallocManaged(&dev_ptr, 128, hipMemAttachGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocManaged: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipHostFree +static bool NoGpuTst_hipHostFree() { + bool passed = false; + hipError_t err; + void* ptr1 = nullptr; + err = hipHostFree(ptr1); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostFree: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipMemcpyWithStream +static bool NoGpuTst_hipMemcpyWithStream() { + bool passed = false; + hipError_t err; + int *A_h, *B_h; + A_h = reinterpret_cast(malloc(sizeof(int))); + REQUIRE(A_h != nullptr); + B_h = reinterpret_cast(malloc(sizeof(int))); + REQUIRE(B_h != nullptr); + err = hipMemcpyWithStream(A_h, B_h, 128, hipMemcpyHostToHost, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMemcpyWithStream: " << + hipGetErrorName(err)); + } + return passed; +} + +#if HT_AMD +// No GPU test case for hipMallocMipmappedArray +static bool NoGpuTst_hipMallocMipmappedArray() { + bool passed = false; + hipError_t err; + hipMipmappedArray_t mipmappedArray{}; + hipExtent extent; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + err = hipMallocMipmappedArray(&mipmappedArray, &desc, extent, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocMipmappedArray: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipDeviceEnablePeerAccess +static bool NoGpuTst_hipDeviceEnablePeerAccess() { + bool passed = false; + hipError_t err; + int peerDeviceId = 0; + err = hipDeviceEnablePeerAccess(peerDeviceId, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceEnablePeerAccess: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipHostMalloc() API +static bool NoGpuTst_hipHostMalloc() { + bool passed = false; + hipError_t err; + int *ptr; + size_t size = 1024; + unsigned int flag = hipHostMallocDefault; + err = hipHostMalloc(&ptr, size, flag); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostMalloc: " <(malloc(size)); + err = hipHostRegister(hostPtr, size, reg_flag); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostRegister: " <(malloc(size)); + int *src = reinterpret_cast(malloc(size)); + err = hipMemcpy(dst, src, size, hipMemcpyHostToHost); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMemcpy: " <(); + hipExtent extent = make_hipExtent(64, 64, 64); + err = hipMalloc3DArray(&array, &desc, extent, hipArrayDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMalloc3DArray: " <(&fn), attr, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetAttribute: " + <(&fn), config); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetCacheConfig: " + <(&fn)); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncGetAttributes: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipModuleLoadData() API +static bool NoGpuTst_hipModuleLoadData() { + bool passed = false; + hipError_t err; + hipModule_t module; + std::ifstream file("dummy_kernel.code", std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) { + err = hipModuleLoadData(&module, &buffer[0]); +#if HT_AMD + if (err == hipErrorNoDevice) { + passed = true; + } +#else + if (err == hipErrorNotInitialized) { + passed = true; + } +#endif + if (passed == false) { + WARN("Error Code Returned by hipModuleLoadData: " << + hipGetErrorName(err)); + } + } else { + WARN("File Read Failed"); + } + return passed; +} + +// No GPU test for hipModuleLoadDataEx() API +static bool NoGpuTst_hipModuleLoadDataEx() { + bool passed = false; + hipError_t err; + hipModule_t module; + std::ifstream file("dummy_kernel.code", std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) { + err = hipModuleLoadDataEx(&module, &buffer[0], 0, nullptr, nullptr); +#if HT_AMD + if (err == hipErrorNoDevice) { + passed = true; + } +#else + if (err == hipErrorNotInitialized) { + passed = true; + } +#endif + if (passed == false) { + WARN("Error Code Returned by hipModuleLoadDataEx: " << + hipGetErrorName(err)); + } + } else { + WARN("File Read Failed"); + } + return true; +} + +// No GPU test for hipLaunchCooperativeKernel() API +static bool NoGpuTst_hipLaunchCooperativeKernel() { + bool passed = false; + hipError_t err; + dim3 dimBlock = dim3(1); + dim3 dimGrid = dim3(1); + err = hipLaunchCooperativeKernel(reinterpret_cast(fn), + dimGrid, dimBlock, nullptr, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipLaunchCooperativeKernel: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipLaunchCooperativeKernelMultiDevice() API +static bool NoGpuTst_hipLaunchCooperativeKernelMultiDevice() { + bool passed = false; + hipError_t err; + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + hipLaunchParams md_params[1]; + void *dev_params[1][1]; + dev_params[0][0] = nullptr; + md_params[0].func = reinterpret_cast(fn); + md_params[0].gridDim = dimGrid; + md_params[0].blockDim = dimBlock; + md_params[0].sharedMem = 0; + md_params[0].stream = hipStreamPerThread; + md_params[0].args = dev_params[0]; + + err = hipLaunchCooperativeKernelMultiDevice(md_params, 1, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipLaunchCooperativeKernelMultiDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipExtLaunchMultiKernelMultiDevice() API +#if HT_AMD +static bool NoGpuTst_hipExtLaunchMultiKernelMultiDevice() { + bool passed = false; + hipError_t err; + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + hipLaunchParams md_params[1]; + md_params[0].func = reinterpret_cast(fn); + md_params[0].gridDim = dimGrid; + md_params[0].blockDim = dimBlock; + md_params[0].sharedMem = 0; + md_params[0].stream = hipStreamPerThread; + md_params[0].args = nullptr; + + err = hipExtLaunchMultiKernelMultiDevice(md_params, 1, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtLaunchMultiKernelMultiDevice: " << + hipGetErrorName(err)); + } + return passed; +} +#endif + +// No GPU test for hipExtLaunchMultiKernelMultiDevice() API +static bool NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor() { + bool passed = false; + hipError_t err; + int numBlock = 0, blockSize = 256; + err = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, + fn, blockSize, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code hipOccupancyMaxActiveBlocksPerMultiprocessor: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags() +static bool NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags() { + bool passed = false; + hipError_t err; + int numBlock = 0, blockSize = 256; + err = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(&numBlock, + fn, blockSize, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags: " + << hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipOccupancyMaxPotentialBlockSize() +static bool NoGpuTst_hipOccupancyMaxPotentialBlockSize() { + bool passed = false; + hipError_t err; + int blocksizelimit = 256, blockSize = 0, gridsize = 0; + err = hipOccupancyMaxPotentialBlockSize(&gridsize, &blockSize, fn, + 0, blocksizelimit); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipOccupancyMaxPotentialBlockSize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hiprtcVersion() +static bool NoGpuTst_hiprtcVersion() { + bool passed = false; + int major = 0, minor = 0; + hiprtcResult err = hiprtcVersion(&major, &minor); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + WARN("Error Code Returned by hiprtcVersion: " << + hiprtcGetErrorString(err)); + } + return passed; +} + +// No GPU test for hiprtcCreateProgram() and hiprtcDestroyProgram() +static bool NoGpuTst_hiprtcCreateProgram_DestroyProg() { + bool passed = false; + hiprtcProgram prog; + hiprtcResult err = hiprtcCreateProgram(&prog, rtcfn, "rtcfn.cu", + 0, nullptr, nullptr); + if (err == HIPRTC_SUCCESS) { + passed = true; + err = hiprtcDestroyProgram(&prog); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + passed = false; + WARN("Error Code Returned by hiprtcDestroyProgram: " << + hiprtcGetErrorString(err)); + } + } else { + WARN("Error Code Returned by hiprtcCreateProgram: " << + hiprtcGetErrorString(err)); + } + return passed; +} + +// No GPU test for hiprtcAddNameExpression() +static bool NoGpuTst_hiprtcAddNameExpression() { + bool passed = false; + hiprtcProgram prog; + hiprtcResult err = hiprtcCreateProgram(&prog, rtcfn, "rtcfn.cu", + 0, nullptr, nullptr); + if (err == HIPRTC_SUCCESS) { + err = hiprtcAddNameExpression(prog, "dummy_func"); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + passed = false; + WARN("Error Code Returned by hiprtcAddNameExpression: " << + hiprtcGetErrorString(err)); + } + } + if (err == HIPRTC_SUCCESS) { + err = hiprtcDestroyProgram(&prog); + } + return passed; +} + +// No GPU test for hipMallocPitch() +#define SIZE_H 20 +#define SIZE_W 179 +static bool NoGpuTst_hipMallocPitch() { + bool passed = false; + hipError_t err; + int* devPtr; + size_t devPitch; + err = hipMallocPitch(reinterpret_cast(&devPtr), &devPitch, + SIZE_W*sizeof(int), SIZE_H); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocPitch: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipMallocArray() +static bool NoGpuTst_hipMallocArray() { + bool passed = false; + hipError_t err; + hipChannelFormatDesc channelDesc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray* hipArray; + err = hipMallocArray(&hipArray, &channelDesc, 256, 256); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocArray: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test for hipMipmappedArrayCreate() +static bool NoGpuTst_hipMipmappedArrayCreate() { + bool passed = false; + hipError_t err; + unsigned int width = 256, height = 256, mipmap_level = 2; + unsigned int orig_width = width; + unsigned int orig_height = height; + width /= pow(2, mipmap_level); + height /= pow(2, mipmap_level); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, + 0, hipChannelFormatKindFloat); + HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc; + memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR)); + mipmapped_array_desc.Width = orig_width; + mipmapped_array_desc.Height = orig_height; + mipmapped_array_desc.Depth = 0; + mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT; + mipmapped_array_desc.NumChannels = + ((channelDesc.x != 0) + (channelDesc.y != 0) +\ + (channelDesc.z != 0) + (channelDesc.w != 0)); + mipmapped_array_desc.Flags = 0; + + hipMipmappedArray* mip_array_ptr; + err = hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc, + 2*mipmap_level); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMipmappedArrayCreate: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +#define SMALL_SIZE 4 +static bool NoGpuTst_hipMalloc3D() { + bool passed = false; + hipError_t err; + size_t width = SMALL_SIZE * sizeof(char); + size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; + hipPitchedPtr devPitchedPtr; + hipExtent extent = make_hipExtent(width, height, depth); + err = hipMalloc3D(&devPitchedPtr, extent); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMalloc3D: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipGraphCreate() API +static bool NoGpuTst_hipGraphCreate() { + bool passed = false; + hipError_t err; + hipGraph_t graph; + + err = hipGraphCreate(&graph, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGraphCreate: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipStreamBeginCapture() API +static bool NoGpuTst_hipStreamBeginCapture() { + bool passed = false; + hipError_t err; + + err = hipStreamBeginCapture(0, hipStreamCaptureModeGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamBeginCapture with null stream: " << + hipGetErrorName(err)); + } + + err = hipStreamBeginCapture(hipStreamPerThread, hipStreamCaptureModeGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamBeginCapture with " + "hipStreamPerThread stream: " << hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipStreamIsCapturing() API +static bool NoGpuTst_hipStreamIsCapturing() { + bool passed = false; + hipError_t err; + hipStreamCaptureStatus cStatus; + + err = hipStreamIsCapturing(0, &cStatus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamIsCapturing: " << + hipGetErrorName(err)); + } + err = hipStreamIsCapturing(hipStreamPerThread, &cStatus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamIsCapturing with " + "hipStreamPerThread stream: " << hipGetErrorName(err)); + } + return passed; +} +// Common function invoking individual tests +static bool NoGpuTst_Common(bool(*test_fn)()) { + bool passed = false; + int *Ptr = reinterpret_cast(mmap(NULL, sizeof(int), + PROT_READ | PROT_WRITE, + MAP_SHARED|MAP_ANONYMOUS, 0, 0)); + if (Ptr == MAP_FAILED) { + WARN("Mapping Failed\n"); + return false; + } + // forking to get number of gpus on this system + if (fork() == 0) { // child process + int TotlGpus = 0; + static_cast(hipGetDeviceCount(&TotlGpus)); + Ptr[0] = TotlGpus; + exit(0); + } else { // parent process + wait(NULL); + if (Ptr[0] > 0) { // If there are gpus active on the system + if ((setenv(visibility_env_var, "-1", 1)) != 0) { + WARN("Unable to turn on HIP_VISIBLE_DEVICES, hence exiting!"); + return false; + } + passed = test_fn(); + } else if (Ptr[0] == 0) { + WARN("GPUs not available!"); + passed = test_fn(); + } + } + return passed; +} + +// Tests +TEST_CASE("Unit_NoGpuTst_hipGetDeviceCount") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceCount)); +} + +TEST_CASE("Unit_NoGpuTst_hipSetDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipSetDevice)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetP2PAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetP2PAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetPCIBusId") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetPCIBusId)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetByPCIBusId") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetByPCIBusId)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceTotalMem") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceTotalMem)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncSetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreateWithPriority") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreateWithPriority)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamSynchronize") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamSynchronize)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamGetFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamGetFlags)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtStreamCreateWithCUMask") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtStreamCreateWithCUMask)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipStreamAddCallback") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamAddCallback)); +} + +TEST_CASE("Unit_NoGpuTst_hipEventCreateWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipEventCreateWithFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSynchronize") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSynchronize)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDevice)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipPointerGetAttributes") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipPointerGetAttributes)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtMallocWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtMallocWithFlags)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipMallocManaged") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocManaged)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostFree") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostFree)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemcpyWithStream") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemcpyWithStream)); +} + +TEST_CASE("Unit_NoGpuTst_hipMallocArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocArray)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipMallocMipmappedArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocMipmappedArray)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipDeviceEnablePeerAccess") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceEnablePeerAccess)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostMalloc") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostMalloc)); +} + +TEST_CASE("Unit_NoGpuTst_hipMalloc") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostRegister") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostRegister)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemcpy") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemcpy)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemAllocPitch") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemAllocPitch)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemGetInfo") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemGetInfo)); +} + +TEST_CASE("Unit_NoGpuTst_hipMalloc3DArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc3DArray)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceCanAccessPeer") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceCanAccessPeer)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceDisablePeerAccess") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceDisablePeerAccess)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGet") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGet)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceComputeCapability") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceComputeCapability)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetName") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetName)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDeviceProperties") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceProperties)); +} + +TEST_CASE("Unit_NoGpuTst_hipChooseDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipChooseDevice)); +} + +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtGetLinkTypeAndHopCount") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtGetLinkTypeAndHopCount)); +} + +TEST_CASE("Unit_NoGpuTst_hipExtStreamGetCUMask") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtStreamGetCUMask)); +} +#endif + +TEST_CASE("Unit_NoGpuTst_hipFuncSetAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncSetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreateWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreateWithFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetStreamPriorityRange") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetStreamPriorityRange)); +} + +TEST_CASE("Unit_NoGpuTst_hipEventCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipEventCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamGetPriority") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamGetPriority)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceReset") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceReset)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetLimit") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetLimit)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDeviceFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipSetDeviceFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipSetDeviceFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoad") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoad)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncGetAttributes") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncGetAttributes)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoadData") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoadData)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoadDataEx") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoadDataEx)); +} + +TEST_CASE("Unit_NoGpuTst_hipLaunchCooperativeKernel") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipLaunchCooperativeKernel)); +} + +TEST_CASE("Unit_NoGpuTst_hipLaunchCooperativeKernelMultiDevice") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipLaunchCooperativeKernelMultiDevice)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtLaunchMultiKernelMultiDevice") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipExtLaunchMultiKernelMultiDevice)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor)); +} + +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxPotentialBlockSize") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxPotentialBlockSize)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcVersion") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcVersion)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcCreateProgram_DestroyProg") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcCreateProgram_DestroyProg)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcAddNameExpression") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcAddNameExpression)); +} + +TEST_CASE("Unit_NoGpuTst_hipMallocPitch") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocPitch)); +} + +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipMipmappedArrayCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMipmappedArrayCreate)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipMalloc3D") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc3D)); +} + +TEST_CASE("Unit_NoGpuTst_hipGraphCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGraphCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamBeginCapture") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamBeginCapture)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamIsCapturing") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamIsCapturing)); +} +#endif // #ifdef __linux__ diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index f24c63ad8c..54e7708556 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -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) diff --git a/projects/hip/tests/catch/unit/memory/MemUtils.hh b/projects/hip/tests/catch/unit/memory/MemUtils.hh index a05588c3b7..b80343d004 100644 --- a/projects/hip/tests/catch/unit/memory/MemUtils.hh +++ b/projects/hip/tests/catch/unit/memory/MemUtils.hh @@ -344,10 +344,10 @@ static inline void memsetCheck(T* aPtr, size_t value, memType memType, MultiDDat template 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)); diff --git a/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc b/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc index b70c8f74b1..0b524655a9 100644 --- a/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc +++ b/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc @@ -501,8 +501,8 @@ void DrvMemcpy3D::HostArray_DrvMemcpy3D(bool device_context_change) { /* DeAllocating the memory */ template void DrvMemcpy3D::DeAllocateMemory() { - hipArrayDestroy(arr); - hipArrayDestroy(arr1); + HIP_CHECK(hipArrayDestroy(arr)); + HIP_CHECK(hipArrayDestroy(arr1)); free(hData); } diff --git a/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc b/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc index 971d6cf8c6..774ec6a6fd 100644 --- a/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc +++ b/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc @@ -93,7 +93,7 @@ TEST_CASE("Unit_hipMallocManaged_HostDeviceConcurrent") { std::thread host_thread(HostKernelDouble, Hmm, hPtr, N); KernelDouble<<>>(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]); diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy.cc index 20860d2843..84af63bea0 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy.cc @@ -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 " <(&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(X_d), diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc index 7d9a8e1ccc..27e011089c 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc @@ -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)); diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc index 10924687df..76b20ac574 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc @@ -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)); diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc index 0e6a7aacde..1972a5f10a 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc @@ -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}; diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index 6603243585..f7c8e75b71 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -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}; diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc index 7ea39d30be..bb54701af8 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc @@ -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, diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc index 503763aed8..4c4e08ec2d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -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(&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(X_d), diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc index a8ebe9534f..48e881895d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc @@ -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}; diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc index 1c738f1c76..56a76cbad2 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc @@ -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}; diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc index 1dad378a6c..d9c08f3181 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc @@ -105,7 +105,7 @@ void memcpyTests::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::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::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::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]), diff --git a/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc b/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc index 50747e275d..75c6c87b4e 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc @@ -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(hostRes)); + HIP_CHECK(hipHostFree(reinterpret_cast(hostRes))); } diff --git a/projects/hip/tests/catch/unit/memory/hipMemset.cc b/projects/hip/tests/catch/unit/memory/hipMemset.cc index 600a008436..c4adb15e53 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemset.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemset.cc @@ -260,21 +260,21 @@ TEST_CASE("Unit_hipMemset_2AsyncOperations") { std::vector v; v.resize(2048); float* p2, *p3; - hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2); + HIP_CHECK(hipMalloc(reinterpret_cast(&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); diff --git a/projects/hip/tests/catch/unit/memory/hipMemset2D.cc b/projects/hip/tests/catch/unit/memory/hipMemset2D.cc index e48b8931f0..2b2cb0463d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemset2D.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemset2D.cc @@ -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); } diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc b/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc index d6f152cec9..2f2b478ae3 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -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 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 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, diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc index 82a0c2879d..929ff2e3c6 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc @@ -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") { diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc b/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc index 6770a6e16c..2a55a2a0a3 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc @@ -379,10 +379,10 @@ void memsetCheck(T* aPtr, size_t value, memSetType memsetType, MultiDData& data, template 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)); diff --git a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc index 8157c8811d..d4ba5960ce 100644 --- a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc +++ b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc @@ -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(-1); attribs->devicePointer = reinterpret_cast(-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(&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(&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}; diff --git a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp index cd9683ea5b..9666dd4693 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp @@ -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);