From 1cd1b3eae29401d0d0428bafb652bc79ce9c86ac Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Wed, 4 Dec 2019 14:06:53 -0500
Subject: [PATCH] P4 to Git Change 2040884 by jujiang@JJ-HIP on 2019/12/04
 13:59:15

	SWDEV-214490 - Update HIP RT for texture3D in HIP/PAL on Windows
	-Update hipMemcpy3D function
	http://ocltc.amd.com/reviews/r/18346/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#90 edit
---
 hipamd/api/hip/hip_memory.cpp | 29 +++++++++++------------------
 1 file changed, 11 insertions(+), 18 deletions(-)

diff --git a/hipamd/api/hip/hip_memory.cpp b/hipamd/api/hip/hip_memory.cpp
index bd82e2634d..89b06aca95 100644
--- a/hipamd/api/hip/hip_memory.cpp
+++ b/hipamd/api/hip/hip_memory.cpp
@@ -1117,35 +1117,28 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
   size_t srcOrigin[3];
   size_t dstOrigin[3];
   size_t region[3];
+
+  region[2] = p->extent.depth;
+  region[1] = p->extent.height;
+  region[0] = p->extent.width;
+  srcOrigin[0] = p->srcPos.x;
+  srcOrigin[1] = p->srcPos.y;
+  srcOrigin[2] = p->srcPos.z;
+  dstOrigin[0] = p->dstPos.x;
+  dstOrigin[1] = p->dstPos.y;
+  dstOrigin[2] = p->dstPos.z;
+
   if (p->dstArray != nullptr) {
     getByteSizeFromChannelFormatKind(p->dstArray->desc.f, &byteSize);
-    region[2] = p->extent.depth;
-    region[1] = p->extent.height;
-    region[0] = p->extent.width;
-    srcOrigin[0] = p->srcPos.x;
-    srcOrigin[1] = p->srcPos.y;
-    srcOrigin[2] = p->srcPos.z;
     dstPitchInbytes = p->dstArray->width * byteSize;
     srcPitchInBytes = p->srcPtr.pitch;
     srcPtr = (void*)p->srcPtr.ptr;
     dstPtr = p->dstArray->data;
-    dstOrigin[0] = p->dstPos.x;
-    dstOrigin[1] = p->dstPos.y;
-    dstOrigin[2] = p->dstPos.z;
   } else {
-    region[2] = p->extent.depth;
-    region[1] = p->extent.height;
-    region[0] = p->extent.width;
-    srcOrigin[0] = p->srcXInBytes;
-    srcOrigin[1] = p->srcY;
-    srcOrigin[2] = p->srcZ;
     srcPitchInBytes = p->srcPtr.pitch;
     dstPitchInbytes = p->dstPtr.pitch;
     srcPtr = p->srcPtr.ptr;
     dstPtr = p->dstPtr.ptr;
-    dstOrigin[0] = p->dstXInBytes;
-    dstOrigin[1] = p->dstY;
-    dstOrigin[2] = p->dstZ;
   }
 
   // Create buffer rectangle info structure