From b0a1b4bf2aecb35740bbb372fae8991c1fc0b9c2 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Tue, 15 Oct 2019 16:43:10 -0500
Subject: [PATCH] P4 to Git Change 2013937 by kjayapra@1_HIPWS_LNX1_PAL on
 2019/10/15 17:40:44

	SWDEV-203394 - Setting Array Channel Desc in hipArrayCreate.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#77 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#22 edit


[ROCm/clr commit: c6c16b8ebd70197571691040458a439cbfb7e266]
---
 projects/clr/hipamd/api/hip/hip_memory.cpp  |  4 ++
 projects/clr/hipamd/api/hip/hip_texture.cpp | 49 ++++++++++++++++++++-
 2 files changed, 51 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp
index 750deef6e5..f72531c70f 100644
--- a/projects/clr/hipamd/api/hip/hip_memory.cpp
+++ b/projects/clr/hipamd/api/hip/hip_memory.cpp
@@ -36,6 +36,8 @@ extern void getDrvChannelOrderAndType(const enum hipArray_Format Format,
                                       cl_channel_order* channelOrder,
                                       cl_channel_type* channelType);
 
+extern void setDescFromChannelType(cl_channel_type channelType, hipChannelFormatDesc* desc);
+
 amd::Memory* getMemoryObject(const void* ptr, size_t& offset) {
   amd::Memory *memObj = amd::MemObjMap::FindMemObj(ptr);
   if (memObj != nullptr) {
@@ -468,6 +470,8 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat
                             &channelOrder, &channelType);
 
   const cl_image_format image_format = { channelOrder, channelType };
+  setDescFromChannelType(channelType, &(array[0]->desc));
+
   size_t pitch = 0;
   hipError_t status = ihipMallocPitch(ptr, &pitch, array[0]->width, array[0]->height, 1, CL_MEM_OBJECT_IMAGE2D,
                       &image_format);
diff --git a/projects/clr/hipamd/api/hip/hip_texture.cpp b/projects/clr/hipamd/api/hip/hip_texture.cpp
index 28328932c3..b7e6293b95 100644
--- a/projects/clr/hipamd/api/hip/hip_texture.cpp
+++ b/projects/clr/hipamd/api/hip/hip_texture.cpp
@@ -76,6 +76,51 @@ void getDrvChannelOrderAndType(const enum hipArray_Format Format, unsigned int N
   }
 }
 
+void setDescFromChannelType(cl_channel_type channelType, hipChannelFormatDesc* desc) {
+
+  memset(desc, 0x00, sizeof(hipChannelFormatDesc));
+
+  switch (channelType) {
+    case CL_SIGNED_INT8:
+    case CL_SIGNED_INT16:
+    case CL_SIGNED_INT32:
+      desc->f = hipChannelFormatKindSigned;
+      break;
+    case CL_UNSIGNED_INT8:
+    case CL_UNSIGNED_INT16:
+    case CL_UNSIGNED_INT32:
+      desc->f = hipChannelFormatKindUnsigned;
+      break;
+    case CL_HALF_FLOAT:
+    case CL_FLOAT:
+      desc->f = hipChannelFormatKindFloat;
+      break;
+    default:
+      desc->f = hipChannelFormatKindNone;
+      break;
+  }
+
+  switch (channelType) {
+    case CL_SIGNED_INT8:
+    case CL_UNSIGNED_INT8:
+      desc->x = 8;
+      break;
+    case CL_SIGNED_INT16:
+    case CL_UNSIGNED_INT16:
+    case CL_HALF_FLOAT:
+      desc->x = 16;
+      break;
+    case CL_SIGNED_INT32:
+    case CL_UNSIGNED_INT32:
+    case CL_FLOAT:
+      desc->x = 32;
+      break;
+    default:
+      desc->x = 0;
+      break;
+  }
+}
+
 void getChannelOrderAndType(const hipChannelFormatDesc& desc, enum hipTextureReadMode readMode,
                             cl_channel_order* channelOrder, cl_channel_type* channelType) {
     if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w != 0) {
@@ -621,8 +666,8 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi
     HIP_RETURN(hipErrorInvalidImage);
   }
 
-  HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE1D, &offset, tex, array->data, &array->desc, array->width,
-                             array->height, array->depth));
+  HIP_RETURN(ihipBindTexture(CL_MEM_OBJECT_IMAGE2D, &offset, tex, array->data, &array->desc, array->width,
+                             array->height, 0));
 }
 
 hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) {