From 438e7b78b45b1c297d8204e01ff4ea9f75c14f0c Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <sarbojit.sarkar@amd.com>
Date: Fri, 11 Sep 2020 04:34:27 -0400
Subject: [PATCH] Updated hip_porting_guide.md

Change-Id: Iaf32033597513aa2cbfde267487af35317ac67fe


[ROCm/clr commit: b7dc406ccd0c2272f6622914c2af9a443c020f95]
---
 .../hipamd/docs/markdown/hip_porting_guide.md    | 16 ++++++++++++++++
 1 file changed, 16 insertions(+)

diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
index 23e5058061..451eadbe00 100644
--- a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
+++ b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
@@ -496,6 +496,22 @@ int main()
 }
 ```
 
+## CU_POINTER_ATTRIBUTE_MEMORY_TYPE
+To get pointer's memory type in HIP/HIP-Clang one 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.
+
+For example:
+```
+double * ptr;
+hipMalloc(reinterpret_cast<void**>(&ptr), sizeof(double));
+hipPointerAttribute_t attr;
+hipPointerGetAttributes(&attr, ptr); /*attr.memoryType will have value as hipMemoryTypeDevice*/
+
+double* ptrHost;
+hipHostMalloc(&ptrHost, sizeof(double));
+hipPointerAttribute_t attr;
+hipPointerGetAttributes(&attr, ptrHost); /*attr.memoryType will have value as hipMemoryTypeHost*/
+```
+
 ## threadfence_system
 Threadfence_system makes all device memory writes, all writes to mapped host memory, and all writes to peer memory visible to CPU and other GPU devices.
 Some implementations can provide this behavior by flushing the GPU L2 cache.