From eaef4a97b5f8972eaec275ed872e7487d6acde01 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Wed, 11 Dec 2019 16:30:48 -0500
Subject: [PATCH] P4 to Git Change 2044223 by jujiang@JJ-HIP on 2019/12/11
 16:16:37

	SWDEV-214490 - Update HIP RT for texture3D in HIP/PAL on Windows
	-Update hipCreateTextureObject

	http://ocltc.amd.com/reviews/r/18382/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#91 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#30 edit
---
 api/hip/hip_memory.cpp  |  4 ++--
 api/hip/hip_texture.cpp | 22 +++++++++++++++++++---
 2 files changed, 21 insertions(+), 5 deletions(-)

diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp
index 89b06aca95..e6f3f2f622 100644
--- a/api/hip/hip_memory.cpp
+++ b/api/hip/hip_memory.cpp
@@ -500,13 +500,13 @@ hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc
   const cl_image_format image_format = { channelOrder, channelType };
 
  // Dummy flags check
-  switch (flags) {
+   switch (flags) {
     case hipArrayLayered:
+    case hipArrayCubemap:
     case hipArraySurfaceLoadStore:
     case hipArrayTextureGather:
         assert(0 && "Unspported");
         break;
-    case hipArrayCubemap:
     case hipArrayDefault:
     default:
         break;
diff --git a/api/hip/hip_texture.cpp b/api/hip/hip_texture.cpp
index d1c228707b..a4712de1cb 100644
--- a/api/hip/hip_texture.cpp
+++ b/api/hip/hip_texture.cpp
@@ -299,6 +299,8 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
 
   amd::Memory* memory = nullptr;
   size_t offset = 0;
+  cl_mem_object_type clType;
+
   switch (pResDesc->resType) {
     case hipResourceTypeArray:
       {
@@ -316,9 +318,23 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
           case hipArrayTextureGather:
           case hipArrayDefault:
           default:
-            image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
-              CL_MEM_OBJECT_IMAGE2D, memory->getMemFlags(), imageFormat,
-              pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0);
+            switch(pResDesc->res.array.array->textureType) {
+              case hipTextureType3D:
+                clType = CL_MEM_OBJECT_IMAGE3D;
+                image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+                  clType, memory->getMemFlags(), imageFormat,
+                  pResDesc->res.array.array->width, pResDesc->res.array.array->height,
+                  pResDesc->res.array.array->depth, 0, 0);
+                break;
+              case hipTextureType2D:
+                clType = CL_MEM_OBJECT_IMAGE2D;
+                image = new (*hip::getCurrentContext()) amd::Image(*memory->asBuffer(),
+                  clType, memory->getMemFlags(), imageFormat,
+                  pResDesc->res.array.array->width, pResDesc->res.array.array->height, 1, 0, 0);
+                break;
+              default:
+                break;
+            }
             break;
         }
       }