SWDEV-368553 – Add missing mipmap Apis

Add missing mipmap Apis’ implementation.
Fix some bugs of mimpmap apis.
Use hipmipmappedArray to differentiate cuda
and driver apis on Nvidia.
Change-Id: I6079d9f3b2ddf4e42b9a6f7f3902322cfca02cfd

hipamd/src/hip_memory.cpp

@@ -3948,30 +3948,6 @@ hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy) {
   HIP_RETURN(ihipMemcpyParam3D(&desc, nullptr));
 }
 
-hipError_t hipMallocMipmappedArray(hipMipmappedArray_t *mipmappedArray,
-                                   const hipChannelFormatDesc* desc,
-                                   hipExtent extent,
-                                   unsigned int numLevels,
-                                   unsigned int flags) {
-  HIP_INIT_API(hipMallocMipmappedArray, mipmappedArray, desc, extent, numLevels, flags);
-  CHECK_STREAM_CAPTURE_SUPPORTED();
-  HIP_RETURN(hipErrorNotSupported);
-}
-
-hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) {
-  HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray);
-  CHECK_STREAM_CAPTURE_SUPPORTED();
-  HIP_RETURN(hipErrorNotSupported);
-}
-
-hipError_t hipGetMipmappedArrayLevel(hipArray_t *levelArray,
-                                     hipMipmappedArray_const_t mipmappedArray,
-                                     unsigned int level) {
-  HIP_INIT_API(hipGetMipmappedArrayLevel, levelArray, mipmappedArray, level);
-
-  HIP_RETURN(hipErrorNotSupported);
-}
-
 hipError_t ihipMipmapArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr,
                                  HIP_ARRAY3D_DESCRIPTOR* mipmapped_array_desc_ptr,
                                  unsigned int num_mipmap_levels) {
@@ -4029,6 +4005,7 @@ hipError_t ihipMipmapArrayCreate(hipMipmappedArray_t* mipmapped_array_pptr,
   (*mipmapped_array_pptr)->max_mipmap_level = num_mipmap_levels;
   (*mipmapped_array_pptr)->flags = mipmapped_array_desc_ptr->Flags;
   (*mipmapped_array_pptr)->format = mipmapped_array_desc_ptr->Format;
+  (*mipmapped_array_pptr)->num_channels = mipmapped_array_desc_ptr->NumChannels;
 
   return hipSuccess;
 }
@@ -4100,6 +4077,9 @@ hipError_t ihipMipmappedArrayGetLevel(hipArray_t* level_array_pptr,
   (*level_array_pptr)->isDrv = 0;
   (*level_array_pptr)->textureType = 0;
 
+  amd::ScopedLock lock(hip::hipArraySetLock);
+  hip::hipArraySet.insert(*level_array_pptr);
+
   return hipSuccess;
 }
 
@@ -4127,3 +4107,40 @@ hipError_t hipMipmappedArrayGetLevel(hipArray_t* level_array_pptr,
   HIP_RETURN(ihipMipmappedArrayGetLevel(level_array_pptr, mipmapped_array_ptr, mip_level));
 }
 
+hipError_t hipMallocMipmappedArray(hipMipmappedArray_t *mipmappedArray,
+                                   const hipChannelFormatDesc* desc,
+                                   hipExtent extent,
+                                   unsigned int numLevels,
+                                   unsigned int flags) {
+  HIP_INIT_API(hipMallocMipmappedArray, mipmappedArray, desc, extent, numLevels, flags);
+  if (mipmappedArray == nullptr || desc == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  CHECK_STREAM_CAPTURE_SUPPORTED();
+  HIP_ARRAY3D_DESCRIPTOR allocateArray = {extent.width,
+                                          extent.height,
+                                          extent.depth,
+                                          hip::getArrayFormat(*desc),
+                                          hip::getNumChannels(*desc),
+                                          flags};
+  if(!hip::CheckArrayFormat(*desc)) {
+    return hipErrorInvalidValue;
+  }
+  HIP_RETURN(ihipMipmapArrayCreate(mipmappedArray, &allocateArray, numLevels));
+}
+
+hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray) {
+  HIP_INIT_API(hipFreeMipmappedArray, mipmappedArray);
+  CHECK_STREAM_CAPTURE_SUPPORTED();
+  HIP_RETURN(ihipMipmappedArrayDestroy(mipmappedArray));
+}
+
+hipError_t hipGetMipmappedArrayLevel(hipArray_t *levelArray,
+                                     hipMipmappedArray_const_t mipmappedArray,
+                                     unsigned int level) {
+  HIP_INIT_API(hipGetMipmappedArrayLevel, levelArray, mipmappedArray, level);
+  CHECK_STREAM_CAPTURE_SUPPORTED();
+  HIP_RETURN(ihipMipmappedArrayGetLevel(levelArray,
+                                        const_cast<hipMipmappedArray_t>(mipmappedArray),
+                                        level));
+}

hipamd/src/hip_texture.cpp

@@ -139,25 +139,22 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject,
     return hipErrorInvalidValue;
   }
 
-  // Mipmaps are currently not supported.
-  if (pResDesc->resType == hipResourceTypeMipmappedArray) {
-    return hipErrorNotSupported;
-  }
   // We don't program the max_ansio_ratio field in the the HW sampler SRD.
   if (pTexDesc->maxAnisotropy != 0) {
     return hipErrorNotSupported;
   }
   // We don't program the lod_bias field in the HW sampler SRD.
   if (pTexDesc->mipmapLevelBias != 0) {
+    LogError("mipmapLevelBias not supported!");
     return hipErrorNotSupported;
   }
   // We don't program the min_lod field in the HW sampler SRD.
   if (pTexDesc->minMipmapLevelClamp != 0) {
-    return hipErrorNotSupported;
+    LogInfo("minMipmapLevelClamp ignored!");
   }
   // We don't program the max_lod field in the HW sampler SRD.
   if (pTexDesc->maxMipmapLevelClamp != 0) {
-    return hipErrorNotSupported;
+    LogInfo("maxMipmapLevelClamp ignored!");
   }
 
   // TODO ROCclr assumes all dimensions have the same addressing mode.
@@ -242,12 +239,50 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject,
       if (image == nullptr) {
         return hipErrorInvalidValue;
       }
+    } else if (image->parent()) {
+      image->retain();  // Because it will be released as a view in ihipDestroyTextureObject()
     }
     break;
   }
-  case hipResourceTypeMipmappedArray:
-    return hipErrorInvalidValue;
+  case hipResourceTypeMipmappedArray: {
+    cl_mem memObj = reinterpret_cast<cl_mem>(pResDesc->res.array.array->data);
+    if (!is_valid(memObj)) {
+      return hipErrorInvalidValue;
+    }
+    image = as_amd(memObj)->asImage();
 
+    hipTextureReadMode readMode = pTexDesc->readMode;
+    // 32-bit integer format will not be promoted, regardless of whether or not
+    // this hipTextureDesc::readMode is set hipReadModeNormalizedFloat is specified.
+    if ((pResDesc->res.array.array->Format == HIP_AD_FORMAT_SIGNED_INT32) ||
+        (pResDesc->res.array.array->Format == HIP_AD_FORMAT_UNSIGNED_INT32)) {
+      readMode = hipReadModeElementType;
+    }
+
+    // We need to create an image view if the user requested to use normalized pixel values,
+    // due to already having the image created with a different format.
+    if ((pResViewDesc != nullptr) || (readMode == hipReadModeNormalizedFloat) ||
+        (pTexDesc->sRGB == 1)) {
+      // TODO ROCclr currently right now can only change the format of the image.
+      const cl_channel_order channelOrder = (pResViewDesc != nullptr)
+          ? hip::getCLChannelOrder(hip::getNumChannels(pResViewDesc->format), pTexDesc->sRGB)
+          : hip::getCLChannelOrder(pResDesc->res.mipmap.mipmap->num_channels, pTexDesc->sRGB);
+      const cl_channel_type channelType = (pResViewDesc != nullptr)
+          ? hip::getCLChannelType(hip::getArrayFormat(pResViewDesc->format), readMode)
+          : hip::getCLChannelType(pResDesc->res.mipmap.mipmap->format, readMode);
+      const amd::Image::Format imageFormat(cl_image_format{channelOrder, channelType});
+      if (!imageFormat.isValid()) {
+        return hipErrorInvalidValue;
+      }
+
+      image = image->createView(*hip::getCurrentDevice()->asContext(), imageFormat, nullptr, 0, 0,
+                                true);
+      if (image == nullptr) {
+        return hipErrorInvalidValue;
+      }
+    }
+    break;
+  }
   case hipResourceTypeLinear: {
     const cl_channel_order channelOrder = hip::getCLChannelOrder(hip::getNumChannels(pResDesc->res.linear.desc), pTexDesc->sRGB);
     const cl_channel_type channelType = hip::getCLChannelType(hip::getArrayFormat(pResDesc->res.linear.desc), pTexDesc->readMode);