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


[ROCm/clr commit: 38d4f7f091]

projects/clr/hipamd/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;

projects/clr/hipamd/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;
         }
       }