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.
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index 8a466e75d7..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.
  *
@@ -3817,7 +3823,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 +3896,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.