From 905512e255f434fd0c2acff26f87dbf84f5c4b41 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 9 Feb 2023 07:11:30 +0530 Subject: [PATCH 1/3] SWDEV-306306 - Update information for hipMemGetInfo API (#3155) Change-Id: Ib2b7f81189a82fd47f28cd0056f2997d4ead0e34 --- include/hip/hip_runtime_api.h | 20 ++++++++++++++------ 1 file changed, 14 insertions(+), 6 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 8a466e75d7..041f03bbfe 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -3817,7 +3817,7 @@ hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count); * @brief Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant * byte value value. * - * hipMemsetAsync() is asynchronous with respect to the host, so the call may return before the + * hipMemsetAsync() is asynchronous with respect to the host, so the call may return before the * memset is complete. The operation can optionally be associated to a stream by passing a non-zero * stream argument. If stream is non-zero, the operation may overlap with operations in other * streams. @@ -3890,14 +3890,22 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream __dparm(0)); /** * @brief Query memory info. - * Return snapshot of free memory, and total allocatable memory on the device. * - * Returns in *free a snapshot of the current free memory. - * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue - * @warning On HCC, the free memory only accounts for memory allocated by this process and may be - *optimistic. + * On ROCM, this function gets the actual free memory left on the current device, so supports + * the cases while running multi-workload (such as multiple processes, multiple threads, and + * multiple GPUs). + * + * @warning On Windows, the free memory only accounts for memory allocated by this process and may + * be optimistic. + * + * @param[out] free returns free memory on the current device in bytes + * @param[out] total returns total allocatable memory on the current device in bytes + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue + * **/ hipError_t hipMemGetInfo(size_t* free, size_t* total); + hipError_t hipMemPtrGetInfo(void* ptr, size_t* size); /** * @brief Allocate an array on the device. From f331637ff6350776d57305f9dfcb494e06e7b233 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 9 Feb 2023 07:12:07 +0530 Subject: [PATCH 2/3] SWDEV-379789 - Correct virtual function support information (#3154) Change-Id: Ife63b2f4bb3ecdb7fed3b796bc3cd582c1b5ad20 --- docs/markdown/hip_kernel_language.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 2b4f7e9cc8..5f8e9f8f95 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -862,8 +862,9 @@ Output Constraints are specified by an `"="` prefix as shown above ("=v"). This ## C++ Support The following C++ features are not supported: - Run-time-type information (RTTI) -- Virtual functions - Try/catch +- Virtual functions +Virtual functions are not supported if objects containing virtual function tables are passed between GPU's of different offload arch's, e.g. between gfx906 and gfx1030. Otherwise virtual functions are supported. ## Kernel Compilation hipcc now supports compiling C++/HIP kernels to binary code objects. From 121a88db151342df4b5c5fbdee1747945ef9c6d9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 10 Feb 2023 01:42:55 +0530 Subject: [PATCH 3/3] SWDEV-364604 - Amend flag name (#3157) Change-Id: I8fb458268959ea1ca253c4ede815a54f041831cf --- include/hip/hip_runtime_api.h | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 041f03bbfe..f8fccdae35 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -585,9 +585,15 @@ enum hipLimit_t { /** Event can support IPC. hipEventDisableTiming also must be set.*/ #define hipEventInterprocess 0x4 -/** Disable any system scope releases for events. May improve performance. - * The flag is a no-op on CUDA platforms.*/ -#define hipEventDisableReleaseToSystem 0x20000000 +/** Disable performing a system scope sequentially consistent memory fence when the event + * transitions from recording to recorded. This can be used for events that are only being + * used to measure timing, and do not require the event inspection operations + * (see ::hipEventSynchronize, ::hipEventQuery, and ::hipEventElapsedTime) to synchronize-with + * the work on which the recorded event (see ::hipEventRecord) is waiting. + * On some AMD GPU devices this can improve the accuracy of timing measurements by avoiding the + * cost of cache writeback and invalidation, and the performance impact of those actions on the + * execution of following work. */ +#define hipEventDisableSystemFence 0x20000000 /** Use a device-scope release when recording this event. This flag is useful to obtain more * precise timings of commands between events. The flag is a no-op on CUDA platforms.*/ @@ -2420,7 +2426,7 @@ hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, would not record profiling data and provide best performance if used for synchronization. * #hipEventInterprocess : The event can be used as an interprocess event. hipEventDisableTiming flag also must be set when hipEventInterprocess flag is set. - * #hipEventDisableReleaseToSystem : Disable releasing any cached memory to system scope. This may + * #hipEventDisableSystemFence : Disable acquire and release system scope fence. This may improve performance but device memory may not be visible to the host and other devices if this flag is set. *