diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md index b26915c50e..c6e857a90f 100644 --- a/docs/markdown/hip_debugging.md +++ b/docs/markdown/hip_debugging.md @@ -262,7 +262,7 @@ The following is the summary of the most useful environment variables in HIP. | AMD_SERIALIZE_COPY
Serialize copies. | 0 | 1: Wait for completion before enqueue.
2: Wait for completion after enqueue.
3: Both. | | HIP_HOST_COHERENT
Coherent memory in hipHostMalloc. | 0 | 0: memory is not coherent between host and GPU.
1: memory is coherent with host. | | AMD_DIRECT_DISPATCH
Enable direct kernel dispatch. | 1 | 0: Disable.
1: Enable. | - +| GPU_MAX_HW_QUEUES
The maximum number of hardware queues allocated per device. | 4 | The variable controls how many independent hardware queues HIP runtime can create per process, per device. If application allocates more HIP streams than this number, then HIP runtime will reuse the same hardware queues for the new streams in round robin manner. Please note, this maximum number does not apply to either hardware queues that are created for CU masked HIP streams, or cooperative queue for HIP Cooperative Groups (there is only one single queue per device). | ## General Debugging Tips - 'gdb --args' can be used to conveniently pass the executable and arguments to gdb. diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index c4cc2301a6..2b4f7e9cc8 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -455,9 +455,9 @@ Following is the list of supported integer intrinsics. Note that intrinsics are | unsigned int __popcll ( unsigned long long int x )
Count the number of bits that are set to 1 in a 64 bit integer. | | int __mul24 ( int x, int y )
Multiply two 24bit integers. | | unsigned int __umul24 ( unsigned int x, unsigned int y )
Multiply two 24bit unsigned integers. | -[1] +[1] The HIP-Clang implementation of __ffs() and __ffsll() contains code to add a constant +1 to produce the ffs result format. -For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform, +For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform, HIP-Clang provides __lastbit_u32_u32(unsigned int input) and __lastbit_u32_u64(unsigned long long int input). The index returned by __lastbit_ instructions starts at -1, while for ffs the index starts at 0. @@ -496,6 +496,18 @@ long long int clock64() ``` Returns the value of counter that is incremented every clock cycle on device. Difference in values returned provides the cycles used. +``` +long long int wall_clock64() +``` +Returns wall clock count at a constant frequency on the device, which can be queried via HIP API with hipDeviceAttributeWallClockRate attribute of the device in HIP application code, for example, +``` +int wallClkRate = 0; //in kilohertz +HIPCHECK(hipDeviceGetAttribute(&wallClkRate, hipDeviceAttributeWallClockRate, deviceId)); +``` +Where hipDeviceAttributeWallClockRate is a device attribute. +Note that, wall clock frequency is a per-device attribute. + + ## Atomic Functions Atomic functions execute as read-modify-write operations residing in global or shared memory. No other device or thread can observe or modify the memory location during an atomic operation. If multiple instructions from different devices or threads target the same memory location, the instructions are serialized in an undefined order. diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index 1fc966e898..507a72c502 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -102,9 +102,6 @@ A stronger system-level fence can be specified when the event is created with hi - hipEventReleaseToSystem : Perform a system-scope release operation when the event is recorded.  This will make both Coherent and Non-Coherent host memory visible to other agents in the system, but may involve heavyweight operations such as cache flushing.  Coherent memory will typically use lighter-weight in-kernel synchronization mechanisms such as an atomic operation and thus does not need to use hipEventReleaseToSystem. - hipEventDisableTiming: Events created with this flag would not record profiling data and provide best performance if used for synchronization. -Note, for HIP Events used in kernel dispatch using hipExtLaunchKernelGGL/hipExtLaunchKernel, events passed in the API are not explicitly recorded and should only be used to get elapsed time for that specific launch. -In case events are used across multiple dispatches, for example, start and stop events from different hipExtLaunchKernelGGL/hipExtLaunchKernel calls, they will be treated as invalid unrecorded events, HIP will throw error "hipErrorInvalidHandle" from hipEventElapsedTime. - ### Summary and Recommendations: - Coherent host memory is the default and is the easiest to use since the memory is visible to the CPU at typical synchronization points. This memory allows in-kernel synchronization commands such as threadfence_system to work transparently. diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 0685876aef..b76f4b7037 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1215,6 +1215,102 @@ typedef enum hipGraphInstantiateFlags { 1, ///< Automatically free memory allocated in a graph before relaunching. } hipGraphInstantiateFlags; +/** + * Memory allocation properties + */ +typedef struct hipMemAllocationProp { + hipMemAllocationType type; ///< Memory allocation type + hipMemAllocationHandleType requestedHandleType; ///< Requested handle type + hipMemLocation location; ///< Memory location + void* win32HandleMetaData; ///< Metadata for Win32 handles + struct { + unsigned char compressionType; ///< Compression type + unsigned char gpuDirectRDMACapable; ///< RDMA capable + unsigned short usage; ///< Usage + } allocFlags; +} hipMemAllocationProp; + +/** + * Generic handle for memory allocation + */ +typedef struct ihipMemGenericAllocationHandle* hipMemGenericAllocationHandle_t; + +/** + * @brief Flags for granularity + * @enum + * @ingroup Enumerations + */ +typedef enum hipMemAllocationGranularity_flags { + hipMemAllocationGranularityMinimum = 0x0, ///< Minimum granularity + hipMemAllocationGranularityRecommended = 0x1 ///< Recommended granularity for performance +} hipMemAllocationGranularity_flags; + +/** + * @brief Memory handle type + * @enum + * @ingroup Enumerations + */ +typedef enum hipMemHandleType { + hipMemHandleTypeGeneric = 0x0 ///< Generic handle type +} hipMemHandleType; + +/** + * @brief Memory operation types + * @enum + * @ingroup Enumerations + */ +typedef enum hipMemOperationType { + hipMemOperationTypeMap = 0x1, ///< Map operation + hipMemOperationTypeUnmap = 0x2 ///< Unmap operation +} hipMemOperationType; + +/** + * @brief Subresource types for sparse arrays + * @enum + * @ingroup Enumerations + */ +typedef enum hipArraySparseSubresourceType { + hipArraySparseSubresourceTypeSparseLevel = 0x0, ///< Sparse level + hipArraySparseSubresourceTypeMiptail = 0x1 ///< Miptail +} hipArraySparseSubresourceType; + +/** + * Map info for arrays + */ +typedef struct hipArrayMapInfo { + hipResourceType resourceType; ///< Resource type + union { + hipMipmappedArray mipmap; + hipArray_t array; + } resource; + hipArraySparseSubresourceType subresourceType; ///< Sparse subresource type + union { + struct { + unsigned int level; ///< For mipmapped arrays must be a valid mipmap level. For arrays must be zero + unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero + unsigned int offsetX; ///< X offset in elements + unsigned int offsetY; ///< Y offset in elements + unsigned int offsetZ; ///< Z offset in elements + unsigned int extentWidth; ///< Width in elements + unsigned int extentHeight; ///< Height in elements + unsigned int extentDepth; ///< Depth in elements + } sparseLevel; + struct { + unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero + unsigned long long offset; ///< Offset within mip tail + unsigned long long size; ///< Extent in bytes + } miptail; + } subresource; + hipMemOperationType memOperationType; ///< Memory operation type + hipMemHandleType memHandleType; ///< Memory handle type + union { + hipMemGenericAllocationHandle_t memHandle; + } memHandle; + unsigned long long offset; ///< Offset within the memory + unsigned int deviceBitMask; ///< Device ordinal bit mask + unsigned int flags; ///< flags for future use, must be zero now. + unsigned int reserved[2]; ///< Reserved for future use, must be zero now. +} hipArrayMapInfo; // Doxygen end group GlobalDefs /** @} */ //------------------------------------------------------------------------------------------------- @@ -1253,7 +1349,7 @@ hipError_t hipInit(unsigned int flags); * * @param [out] driverVersion * - * @returns #hipSuccess, #hipErrorInavlidValue + * @returns #hipSuccess, #hipErrorInvalidValue * * @warning The HIP feature set does not correspond to an exact CUDA SDK driver revision. * This function always set *driverVersion to 4 as an approximation though HIP supports @@ -1269,7 +1365,7 @@ hipError_t hipDriverGetVersion(int* driverVersion); * * @param [out] runtimeVersion * - * @returns #hipSuccess, #hipErrorInavlidValue + * @returns #hipSuccess, #hipErrorInvalidValue * * @warning The version definition of HIP runtime is different from CUDA. * On AMD platform, the function returns HIP runtime version, @@ -1284,7 +1380,7 @@ hipError_t hipRuntimeGetVersion(int* runtimeVersion); * @param [out] device * @param [in] ordinal * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceGet(hipDevice_t* device, int ordinal); @@ -1294,7 +1390,7 @@ hipError_t hipDeviceGet(hipDevice_t* device, int ordinal); * @param [out] minor * @param [in] device * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device); /** @@ -1303,7 +1399,7 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device * @param [in] len * @param [in] device * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device); /** @@ -1315,7 +1411,7 @@ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device); * it is still open to changes and may have outstanding issues. * * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorNotInitialized, - * #hipErrorDeInitialized + * #hipErrorDeinitialized */ hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device); /** @@ -1325,7 +1421,7 @@ hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device); * @param [in] srcDevice * @param [in] dstDevice * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr, int srcDevice, int dstDevice); @@ -1335,7 +1431,7 @@ hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr, * @param [in] len * @param [in] device * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device); /** @@ -1343,7 +1439,7 @@ hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device); * @param [out] device handle * @param [in] PCI Bus ID * - * @returns #hipSuccess, #hipErrorInavlidDevice, #hipErrorInvalidValue + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue */ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId); /** @@ -1351,7 +1447,7 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId); * @param [out] bytes * @param [in] device * - * @returns #hipSuccess, #hipErrorInavlidDevice + * @returns #hipSuccess, #hipErrorInvalidDevice */ hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device); // doxygen end initialization @@ -2381,13 +2477,6 @@ hipError_t hipEventSynchronize(hipEvent_t event); * recorded on one or both events (that is, hipEventQuery() would return #hipErrorNotReady on at * least one of the events), then #hipErrorNotReady is returned. * - * Note, for HIP Events used in kernel dispatch using hipExtLaunchKernelGGL/hipExtLaunchKernel, - * events passed in hipExtLaunchKernelGGL/hipExtLaunchKernel are not explicitly recorded and should - * only be used to get elapsed time for that specific launch. In case events are used across - * multiple dispatches, for example, start and stop events from different hipExtLaunchKernelGGL/ - * hipExtLaunchKernel calls, they will be treated as invalid unrecorded events, HIP will throw - * error "hipErrorInvalidHandle" from hipEventElapsedTime. - * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventDestroy, hipEventRecord, * hipEventSynchronize */ @@ -3334,7 +3423,7 @@ hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes, * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -3352,7 +3441,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes); * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -3370,7 +3459,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes); * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -3388,7 +3477,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -3406,7 +3495,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, h * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -3424,7 +3513,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h * @param[in] src Data being copy from * @param[in] sizeBytes Data size in bytes * - * @return #hipSuccess, #hipErrorDeInitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, + * @return #hipSuccess, #hipErrorDeinitialized, #hipErrorNotInitialized, #hipErrorInvalidContext, * #hipErrorInvalidValue * * @see hipArrayCreate, hipArrayDestroy, hipArrayGetDescriptor, hipMemAlloc, hipMemAllocHost, @@ -6400,103 +6489,6 @@ hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, u */ -/** - * Memory allocation properties - */ -typedef struct hipMemAllocationProp { - hipMemAllocationType type; ///< Memory allocation type - hipMemAllocationHandleType requestedHandleType; ///< Requested handle type - hipMemLocation location; ///< Memory location - void* win32HandleMetaData; ///< Metadata for Win32 handles - struct { - unsigned char compressionType; ///< Compression type - unsigned char gpuDirectRDMACapable; ///< RDMA capable - unsigned short usage; ///< Usage - } allocFlags; -} hipMemAllocationProp; - -/** - * Generic handle for memory allocation - */ -typedef struct ihipMemGenericAllocationHandle* hipMemGenericAllocationHandle_t; - -/** - * @brief Flags for granularity - * @enum - * @ingroup Enumerations - */ -typedef enum hipMemAllocationGranularity_flags { - hipMemAllocationGranularityMinimum = 0x0, ///< Minimum granularity - hipMemAllocationGranularityRecommended = 0x1 ///< Recommended granularity for performance -} hipMemAllocationGranularity_flags; - -/** - * @brief Memory handle type - * @enum - * @ingroup Enumerations - */ -typedef enum hipMemHandleType { - hipMemHandleTypeGeneric = 0x0 ///< Generic handle type -} hipMemHandleType; - -/** - * @brief Memory operation types - * @enum - * @ingroup Enumerations - */ -typedef enum hipMemOperationType { - hipMemOperationTypeMap = 0x1, ///< Map operation - hipMemOperationTypeUnmap = 0x2 ///< Unmap operation -} hipMemOperationType; - -/** - * @brief Subresource types for sparse arrays - * @enum - * @ingroup Enumerations - */ -typedef enum hipArraySparseSubresourceType { - hipArraySparseSubresourceTypeSparseLevel = 0x0, ///< Sparse level - hipArraySparseSubresourceTypeMiptail = 0x1 ///< Miptail -} hipArraySparseSubresourceType; - -/** - * Map info for arrays - */ -typedef struct hipArrayMapInfo { - hipResourceType resourceType; ///< Resource type - union { - hipMipmappedArray mipmap; - hipArray_t array; - } resource; - hipArraySparseSubresourceType subresourceType; ///< Sparse subresource type - union { - struct { - unsigned int level; ///< For mipmapped arrays must be a valid mipmap level. For arrays must be zero - unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero - unsigned int offsetX; ///< X offset in elements - unsigned int offsetY; ///< Y offset in elements - unsigned int offsetZ; ///< Z offset in elements - unsigned int extentWidth; ///< Width in elements - unsigned int extentHeight; ///< Height in elements - unsigned int extentDepth; ///< Depth in elements - } sparseLevel; - struct { - unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero - unsigned long long offset; ///< Offset within mip tail - unsigned long long size; ///< Extent in bytes - } miptail; - } subresource; - hipMemOperationType memOperationType; ///< Memory operation type - hipMemHandleType memHandleType; ///< Memory handle type - union { - hipMemGenericAllocationHandle_t memHandle; - } memHandle; - unsigned long long offset; ///< Offset within the memory - unsigned int deviceBitMask; ///< Device ordinal bit mask - unsigned int flags; ///< flags for future use, must be zero now. - unsigned int reserved[2]; ///< Reserved for future use, must be zero now. -} hipArrayMapInfo; - /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- diff --git a/tests/catch/include/hip_test_rtc.hh b/tests/catch/include/hip_test_rtc.hh index 0f862840bb..11ef6a1656 100644 --- a/tests/catch/include/hip_test_rtc.hh +++ b/tests/catch/include/hip_test_rtc.hh @@ -77,11 +77,7 @@ inline std::vector alignArguments(std::vector& args) { for (auto& arg : args) { const char* argPtr{reinterpret_cast(arg.ptr)}; - /* - * Details about the padding formula can be found at: - * https://en.wikipedia.org/wiki/Data_structure_alignment#Data_structure_padding - */ - int paddingNeeded = -count & (arg.alignmentRequirement - 1); + int paddingNeeded = (arg.alignmentRequirement - 1) & (~count + 1); alignedArguments.insert(std::end(alignedArguments), paddingNeeded, 0); count += paddingNeeded; diff --git a/tests/catch/packaging/hip-tests.txt b/tests/catch/packaging/hip-tests.txt index f934130e34..553fdb853b 100644 --- a/tests/catch/packaging/hip-tests.txt +++ b/tests/catch/packaging/hip-tests.txt @@ -71,6 +71,9 @@ set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "HIP Support ") set(CPACK_PACKAGE_VERSION @HIP_VERSION_MAJOR@.@HIP_VERSION_MINOR@.@HIP_VERSION_PATCH_GITHASH@) +# to remove hip-catch-* package during uninstallation of rocm +set (CPACK_DEBIAN_PACKAGE_DEPENDS "rocm-core") +set (CPACK_RPM_PACKAGE_REQUIRES "rocm-core") if(NOT WIN32) set(CPACK_GENERATOR "TGZ;DEB;RPM" CACHE STRING "Linux package types for catch tests") diff --git a/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc b/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc index 11b0faa98f..7911b133ec 100644 --- a/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc +++ b/tests/catch/unit/device/hipExtGetLinkTypeAndHopCount.cc @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include -#if HT_AMD +#if __linux__ && HT_AMD TEST_CASE("Unit_hipExtGetLinkTypeAndHopCount_Positive_Basic") { const auto device1 = GENERATE(range(0, HipTest::getDeviceCount())); const auto device2 = GENERATE(range(0, HipTest::getDeviceCount())); @@ -97,4 +97,4 @@ TEST_CASE("Unit_hipExtGetLinkTypeAndHopCount_Negative_Parameters") { HIP_CHECK_ERROR(hipExtGetLinkTypeAndHopCount(0, 1, nullptr, nullptr), hipErrorInvalidValue); } } -#endif \ No newline at end of file +#endif diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 54e7708556..4d2d74c033 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -97,6 +97,7 @@ set(TEST_SRC hipMemsetSync.cc hipMemsetAsync.cc hipMemAdvise.cc + hipMemRangeGetAttributes.cc ) else() set(TEST_SRC @@ -170,6 +171,7 @@ set(TEST_SRC hipMemsetSync.cc hipMemsetAsync.cc hipMemAdvise.cc + hipMemRangeGetAttributes.cc ) endif() diff --git a/tests/catch/unit/memory/hipMemAdvise.cc b/tests/catch/unit/memory/hipMemAdvise.cc index cc2a52a807..09ea32bda8 100644 --- a/tests/catch/unit/memory/hipMemAdvise.cc +++ b/tests/catch/unit/memory/hipMemAdvise.cc @@ -143,7 +143,7 @@ TEST_CASE("Unit_hipMemAdvise_TstFlags") { float *Hmm = nullptr; int AttrVal = 0; HIP_CHECK(hipGetDeviceCount(&NumDevs)); - Outpt = new int(NumDevs); + Outpt = new int[NumDevs]; HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); // With the following for loop we iterate through each of the Gpus in the // system set and unset the flags and check the behavior. diff --git a/tests/catch/unit/memory/hipMemRangeGetAttributes.cc b/tests/catch/unit/memory/hipMemRangeGetAttributes.cc new file mode 100644 index 0000000000..e1240f4330 --- /dev/null +++ b/tests/catch/unit/memory/hipMemRangeGetAttributes.cc @@ -0,0 +1,325 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* Test Case Description: + Scenario-1: Testing basic working of hipMemRangeGetAttributes() + api with different flags + Scenario-2: Negative testing with hipMemRangeGetAttributes() api +*/ + +#include +#define MEM_SIZE 8192 + +static bool CheckError(hipError_t err, int LineNo) { + if (err == hipSuccess) { + WARN("Error expected but received hipSuccess at line no.:" + << LineNo); + return false; + } else { + return true; + } +} + +static int HmmAttrPrint() { + int managed = 0; + WARN("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" + << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + WARN("hipDeviceAttributeManagedMemory: " << managed); + return managed; +} + +#ifdef __linux__ +/* Test Scenario: Testing basic working of hipMemRangeGetAttributes() + api with different flags */ + +TEST_CASE("Unit_hipMemRangeGetAttributes_TstFlgs") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0; + int *Outpt[4], *AcsdBy = nullptr; + float *Hmm = nullptr; + hipStream_t strm; + hipMemRangeAttribute AttrArr[4] = + {hipMemRangeAttributeReadMostly, + hipMemRangeAttributePreferredLocation, + hipMemRangeAttributeAccessedBy, + hipMemRangeAttributeLastPrefetchLocation}; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + AcsdBy = new int(NumDevs); + size_t dataSizes[4] = {sizeof(int), sizeof(int), + (NumDevs * sizeof(int)), sizeof(int)}; + Outpt[0] = new int; + Outpt[1] = new int; + Outpt[2] = new int[NumDevs]; + Outpt[3] = new int; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + dataSizes, AttrArr, 4, Hmm, + MEM_SIZE)); + if (*(Outpt[0]) != 1) { + WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + + if (*(Outpt[0]) != 0) { + WARN("Attempt to set hipMemAdviseUnsetReadMostly flag failed!\n"); + IfTestPassed = false; + } + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, + hipMemAdviseSetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + if (*(Outpt[1]) != i) { + WARN("Attempt to set hipMemAdviseSetPreferredLocation flag"); + WARN(" failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + if ((Outpt[2][0]) != i) { + WARN("Attempt to set hipMemAdviseSetAccessedBy flag"); + WARN(" failed!\n"); + IfTestPassed = false; + } + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + if (!((Outpt[2][i]) < 0)) { + WARN("Attempt to set hipMemAdviseUnsetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + if (*(Outpt[3]) != i) { + WARN("Attempt to prefetch memory to device: " << i); + WARN("failed!\n"); + IfTestPassed = false; + } + // Prefetching back to Host + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, -1, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE)); + if (*(Outpt[3]) != -1) { + WARN("Attempt to prefetch memory to Host failed!\n"); + IfTestPassed = false; + } + } + + HIP_CHECK(hipFree(Hmm)); + delete[] AcsdBy; + for (int i = 0; i < 4; ++i) { + delete Outpt[i]; + } + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* Test Scenario: Negative testing with hipMemRangeGetAttributes() api*/ +TEST_CASE("Unit_hipMemRangeGetAttributes_NegativeTst") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, *Outpt[4]; + float *Hmm = nullptr; + hipMemRangeAttribute AttrArr[4] = + {hipMemRangeAttributeReadMostly, + hipMemRangeAttributePreferredLocation, + hipMemRangeAttributeAccessedBy, + hipMemRangeAttributeLastPrefetchLocation}; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + size_t dataSizes[4] = {sizeof(int), sizeof(int), + (NumDevs * sizeof(int)), sizeof(int)}; + Outpt[0] = new int; + Outpt[1] = new int; + Outpt[2] = new int[NumDevs]; + Outpt[3] = new int; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + // passing zero for num of attributes param(4th) + SECTION("passing zero for num of attributes param(4th)") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 0, Hmm, MEM_SIZE), __LINE__)) { + IfTestPassed = false; + } + } + + // the first dataSize element passed as 0 + dataSizes[0] = 0; + dataSizes[1] = sizeof(int); + dataSizes[2] = NumDevs * sizeof(int); + dataSizes[3] = sizeof(int); + SECTION("the first dataSize element passed as 0") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // passing datasize as 2 while the requirement is multiple of 4 + dataSizes[0] = 2; + dataSizes[1] = sizeof(int); + dataSizes[2] = NumDevs * sizeof(int); + dataSizes[3] = sizeof(int); + SECTION("datasize as 2 while the requirement is multiple of 4") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // passing datasize as 6 while the requirement is multiple of 4 + dataSizes[0] = 6; + dataSizes[1] = sizeof(int); + dataSizes[2] = NumDevs * sizeof(int); + dataSizes[3] = sizeof(int); + SECTION("datasize as 6 while the requirement is multiple of 4") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // passing datasize as 7 while the requirement is multiple of 4 + dataSizes[0] = 7; + dataSizes[1] = sizeof(int); + dataSizes[2] = NumDevs * sizeof(int); + dataSizes[3] = sizeof(int); + SECTION("datasize as 7 while the requirement is multiple of 4") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // passing dataSize as 7 for attribute hipMemRangeAttributeAccessedBy + hipMemRangeAttribute AttrArr1[1] = {hipMemRangeAttributeAccessedBy}; + dataSizes[2] = {7}; + SECTION("passing dataSize as 7 for attribute hipMemRangeAttrAccessedBy") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr1, 1, Hmm, MEM_SIZE), __LINE__)) { + IfTestPassed = false; + } + } + // Passing NULL as first parameter + SECTION("Passing NULL as first parameter") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(NULL), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // Passing count parameter as zero + SECTION("Passing count parameter as zero") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + AttrArr, 4, Hmm, 0), + __LINE__)) { + IfTestPassed = false; + } + } + // Passing NULL for Attribute array(3rd param) + SECTION("Passing NULL for Attribute array(3rd param)") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + NULL, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + // Passing 0 for Attribute array(3rd param) + SECTION("Passing 0 for Attribute array(3rd param)") { + if (!CheckError(hipMemRangeGetAttributes( + reinterpret_cast(Outpt), + reinterpret_cast(dataSizes), + 0, 4, Hmm, MEM_SIZE), + __LINE__)) { + IfTestPassed = false; + } + } + for (int i = 0; i < 4; ++i) { + delete Outpt[i]; + } + REQUIRE(IfTestPassed); + + // The following scenarios have been removed considering the nature of the + // api. With Consultation with Maneesh Gupta, the following scenarios + // have been removed. + // passing numAttributes as 4 while the attributes array has only 2 members + // passing numAttributes as 10 while the attributes array has only 2 members + // length of the list of dataSizes less than the number of + // attributes being probed + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} +#endif diff --git a/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp b/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp index 0c332308db..ce882b12d0 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipLaunchCoopMultiKernel.cpp @@ -41,7 +41,6 @@ const static uint BufferSizeInDwords = 256 * 1024 * 1024; const static uint numQueues = 4; const static uint numIter = 100; constexpr uint NumKernelArgs = 4; -constexpr uint MaxGPUs = 8; #include @@ -93,10 +92,7 @@ __global__ void test_gws(uint* buf, uint bufSize, long* tmpBuf, long* result) int main() { float *A, *B; - uint* dA[MaxGPUs]; - long* dB[MaxGPUs]; long* dC; - hipStream_t stream[MaxGPUs]; uint32_t* init = new uint32_t[BufferSizeInDwords]; for (uint32_t i = 0; i < BufferSizeInDwords; ++i) { @@ -106,7 +102,11 @@ int main() { int nGpu = 0; HIPCHECK(hipGetDeviceCount(&nGpu)); size_t copySizeInDwords = BufferSizeInDwords / nGpu; - hipDeviceProp_t deviceProp[MaxGPUs]; + + uint* dA[nGpu]; + long* dB[nGpu]; + hipStream_t stream[nGpu]; + hipDeviceProp_t deviceProp[nGpu]; for (int i = 0; i < nGpu; i++) { HIPCHECK(hipSetDevice(i)); @@ -146,7 +146,7 @@ int main() { std::time_t end_time; double time = 0; for (uint set = 0; set < 3; ++set) { - void* args[MaxGPUs * NumKernelArgs]; + void* args[nGpu * NumKernelArgs]; std::cout << "---------- Test#" << set << ", size: "<< BufferSizeInDwords << " dwords ---------------\n"; for (int i = 0; i < nGpu; i++) { @@ -205,6 +205,7 @@ int main() { hipFree(dB[i]); HIPCHECK(hipStreamDestroy(stream[i])); } + delete [] init; passed(); return 0; diff --git a/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp b/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp index fad2eb8bb7..14d9eab597 100644 --- a/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp +++ b/tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp @@ -136,7 +136,7 @@ int main(int argc, char* argv[]) { // Therefore the hipStreamSynchronize() in the // main thread should hardly take any time to complete. - if (duration.count() < 100) { + if (duration.count() < 200) { passed(); } else { failed("hipStreamSynchronize is taking more time than expected after Callback()");