[SYCLomatic] Add migration support for eight mipmap texture API (#2743)

Signed-off-by: Deepak Raj H R <deepak.raj.h.r@intel.com>

Author: the-slow-one

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -557,6 +557,11 @@ void MapNames::setExplicitNamespaceMap(
            DpctGlobalInfo::useExtBindlessImages()
                ? getDpctNamespace() + "experimental::image_mem_wrapper_ptr"
                : "cudaMipmappedArray_t")},
+      {"CUmipmappedArray",
+       std::make_shared<TypeNameRule>(
+           DpctGlobalInfo::useExtBindlessImages()
+               ? getDpctNamespace() + "experimental::image_mem_wrapper_ptr"
+               : "CUmipmappedArray")},
       {"cudaTextureDesc",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "sampling_info",
                                       HelperFeatureEnum::device_ext)},

clang/lib/DPCT/RulesLang/APINamesMemory.inc

@@ -185,6 +185,14 @@ CONDITIONAL_FACTORY_ENTRY(
     REMOVE_API_FACTORY_ENTRY_WITH_MSG("cudaGetMipmappedArrayLevel",
                                       MipmapNeedBindlessImage))
 
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cuMipmappedArrayGetLevel", DEREF(makeCallArgCreatorWithCall(0)),
+        MEMBER_CALL(ARG(1), true, "get_mip_level", ARG(2)))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuMipmappedArrayGetLevel",
+                                      MipmapNeedBindlessImage))
+
 CONDITIONAL_FACTORY_ENTRY(
     UseExtBindlessImages,
     ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(

clang/lib/DPCT/RulesLang/APINamesTexture.inc

@@ -314,6 +314,67 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                                           "image_matrix",
                                 ARG(1)))))
 ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0)))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cuMipmappedArrayCreate", DEREF(0),
+        NEW(DpctGlobalInfo::useExtBindlessImages()
+                ? MapNames::getDpctNamespace() +
+                      "experimental::image_mem_wrapper"
+                : MapNames::getDpctNamespace() + "image_matrix",
+            DEREF(1), ARG(2)))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuMipmappedArrayCreate",
+                                      MipmapNeedBindlessImage))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuMipmappedArrayDestroy",
+                                             ARG(0))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuMipmappedArrayDestroy",
+                                      MipmapNeedBindlessImage))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(MEMBER_CALL_FACTORY_ENTRY("cuTexRefSetMipmappedArray",
+                                                 ARG(0), true, "attach",
+                                                 ARG(1))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuTexRefSetMipmappedArray",
+                                      MipmapNeedBindlessImage))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(MEMBER_CALL_FACTORY_ENTRY("cuTexRefSetMipmapFilterMode",
+                                                 ARG(0), true,
+                                                 "set_mip_filtering_mode",
+                                                 ARG(1))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuTexRefSetMipmapFilterMode",
+                                      MipmapNeedBindlessImage))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cuTexRefGetMipmapFilterMode", DEREF(0),
+        MEMBER_CALL(ARG(1), true, "get_mip_filtering_mode"))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuTexRefGetMipmapFilterMode",
+                                      MipmapNeedBindlessImage))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(MEMBER_CALL_FACTORY_ENTRY("cuTexRefGetMipmapLevelClamp",
+                                                 ARG(2), true,
+                                                 "get_mip_level_clamp", ARG(0),
+                                                 ARG(1))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuTexRefGetMipmapLevelClamp",
+                                      MipmapNeedBindlessImage))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cuTexRefGetMipmappedArray", DEREF(0),
+        MEMBER_CALL(ARG(1), true, "get_attached_mipmap_data"))),
+    REMOVE_API_FACTORY_ENTRY_WITH_MSG("cuTexRefGetMipmappedArray",
+                                      MipmapNeedBindlessImage))
+
 ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED)
 FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,

clang/lib/DPCT/RulesLang/RulesLangTexture.cpp

@@ -520,7 +520,8 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
               "CUresourcetype", "CUresourcetype_enum", "CUaddress_mode",
               "CUaddress_mode_enum", "CUfilter_mode", "CUfilter_mode_enum",
               "CUDA_TEXTURE_DESC", "CUtexref", "textureReference",
-              "cudaMipmappedArray", "cudaMipmappedArray_t"))))))
+              "cudaMipmappedArray", "cudaMipmappedArray_t",
+              "CUmipmappedArray"))))))
           .bind("texType"),
       this);
 
@@ -574,6 +575,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "cuSurfObjectDestroy",
       "cuArray3DCreate_v2",
       "cuArrayCreate_v2",
+      "cuMipmappedArrayCreate",
+      "cuMipmappedArrayDestroy",
+      "cuMipmappedArrayGetLevel",
       "cuArrayDestroy",
       "cuTexObjectCreate",
       "cuTexObjectDestroy",
@@ -589,6 +593,11 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "cuTexRefGetFlags",
       "cuTexRefSetAddress_v2",
       "cuTexRefSetAddress2D_v3",
+      "cuTexRefSetMipmappedArray",
+      "cuTexRefGetMipmappedArray",
+      "cuTexRefGetMipmapFilterMode",
+      "cuTexRefSetMipmapFilterMode",
+      "cuTexRefGetMipmapLevelClamp",
   };
 
   auto hasAnyFuncName = [&]() {

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -1712,9 +1712,9 @@ ENTRY(cuMemsetD32, cuMemsetD32_v2, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemsetD32Async, cuMemsetD32Async, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemsetD8, cuMemsetD8_v2, true, NO_FLAG, P4, "comment")
 ENTRY(cuMemsetD8Async, cuMemsetD8Async, true, NO_FLAG, P4, "comment")
-ENTRY(cuMipmappedArrayCreate, cuMipmappedArrayCreate, false, NO_FLAG, P4, "comment")
-ENTRY(cuMipmappedArrayDestroy, cuMipmappedArrayDestroy, false, NO_FLAG, P4, "comment")
-ENTRY(cuMipmappedArrayGetLevel, cuMipmappedArrayGetLevel, false, NO_FLAG, P4, "comment")
+ENTRY(cuMipmappedArrayCreate, cuMipmappedArrayCreate, true, NO_FLAG, P4, "Successful")
+ENTRY(cuMipmappedArrayDestroy, cuMipmappedArrayDestroy, true, NO_FLAG, P4, "Successful")
+ENTRY(cuMipmappedArrayGetLevel, cuMipmappedArrayGetLevel, true, NO_FLAG, P4, "Successful")
 ENTRY(cuMipmappedArrayGetMemoryRequirements, cuMipmappedArrayGetMemoryRequirements, false, NO_FLAG, P4, "comment")
 ENTRY(cuMipmappedArrayGetSparseProperties, cuMipmappedArrayGetSparseProperties, false, NO_FLAG, P7, "comment")
 
@@ -1962,10 +1962,10 @@ ENTRY(cuTexRefGetFilterMode, cuTexRefGetFilterMode, true, NO_FLAG, P4, "Successf
 ENTRY(cuTexRefGetFlags, cuTexRefGetFlags, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefGetFormat, cuTexRefGetFormat, false, NO_FLAG, P4, "comment")
 ENTRY(cuTexRefGetMaxAnisotropy, cuTexRefGetMaxAnisotropy, false, NO_FLAG, P4, "comment")
-ENTRY(cuTexRefGetMipmapFilterMode, cuTexRefGetMipmapFilterMode, false, NO_FLAG, P4, "comment")
+ENTRY(cuTexRefGetMipmapFilterMode, cuTexRefGetMipmapFilterMode, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefGetMipmapLevelBias, cuTexRefGetMipmapLevelBias, false, NO_FLAG, P4, "comment")
-ENTRY(cuTexRefGetMipmapLevelClamp, cuTexRefGetMipmapLevelClamp, false, NO_FLAG, P4, "comment")
-ENTRY(cuTexRefGetMipmappedArray, cuTexRefGetMipmappedArray, false, NO_FLAG, P4, "comment")
+ENTRY(cuTexRefGetMipmapLevelClamp, cuTexRefGetMipmapLevelClamp, true, NO_FLAG, P4, "Successful")
+ENTRY(cuTexRefGetMipmappedArray, cuTexRefGetMipmappedArray, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefSetAddress, cuTexRefSetAddress_v2, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefSetAddress2D, cuTexRefSetAddress2D_v2, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefSetAddressMode, cuTexRefSetAddressMode, true, NO_FLAG, P4, "Successful")
@@ -1975,10 +1975,10 @@ ENTRY(cuTexRefSetFilterMode, cuTexRefSetFilterMode, true, NO_FLAG, P4, "Successf
 ENTRY(cuTexRefSetFlags, cuTexRefSetFlags, true, NO_FLAG, P4, "DPCT1074")
 ENTRY(cuTexRefSetFormat, cuTexRefSetFormat, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefSetMaxAnisotropy, cuTexRefSetMaxAnisotropy, false, NO_FLAG, P4, "comment")
-ENTRY(cuTexRefSetMipmapFilterMode, cuTexRefSetMipmapFilterMode, false, NO_FLAG, P4, "comment")
+ENTRY(cuTexRefSetMipmapFilterMode, cuTexRefSetMipmapFilterMode, true, NO_FLAG, P4, "Successful")
 ENTRY(cuTexRefSetMipmapLevelBias, cuTexRefSetMipmapLevelBias, false, NO_FLAG, P4, "comment")
 ENTRY(cuTexRefSetMipmapLevelClamp, cuTexRefSetMipmapLevelClamp, false, NO_FLAG, P4, "comment")
-ENTRY(cuTexRefSetMipmappedArray, cuTexRefSetMipmappedArray, false, NO_FLAG, P4, "comment")
+ENTRY(cuTexRefSetMipmappedArray, cuTexRefSetMipmappedArray, true, NO_FLAG, P4, "Successful")
 
 // Surface Reference Management(Deprecated)
 ENTRY(cuSurfRefGetArray, cuSurfRefGetArray, false, NO_FLAG, P4, "comment")

clang/runtime/dpct-rt/include/dpct/bindless_images.hpp

@@ -420,6 +420,19 @@ class image_mem_wrapper {
 
     init_mip_level_wrappers(q);
   }
+  /// Create bindless image memory wrapper.
+  /// \param [in] desc The image descriptor of bindless image.
+  /// \param [in] num_levels Number of mipmap levels to allocate
+  image_mem_wrapper(sycl::ext::oneapi::experimental::image_descriptor desc,
+                    unsigned int num_levels)
+      : _desc(desc) {
+    _desc.type = sycl::ext::oneapi::experimental::image_type::mipmap;
+    _desc.num_levels = num_levels;
+    auto q = get_default_queue();
+    _handle = alloc_image_mem(_desc, q);
+    init_mip_level_wrappers(q);
+  }
+
   image_mem_wrapper(const image_mem_wrapper &) = delete;
   image_mem_wrapper &operator=(const image_mem_wrapper &) = delete;
   /// Destroy bindless image memory wrapper.
@@ -462,6 +475,10 @@ class image_mem_wrapper {
     return _sub_wrappers + level;
   }
 
+  sycl::ext::oneapi::experimental::image_type get_image_type(void) {
+    return _desc.type;
+  }
+
 private:
   image_mem_wrapper(
       const image_channel &channel,
@@ -564,11 +581,11 @@ class external_mem_wrapper_base {
   image_mem_wrapper *get_mapped_mipmapped_array() {
     if (_res_is_buffer) {
       throw std::runtime_error(
-          "Buffer resouce cannot be accessed as an array!");
+          "Buffer resource cannot be accessed as an array!");
     }
     if (!_res_img_mem_wrapper_ptr) {
       throw std::runtime_error(
-          "Resouce is not mapped! "
+          "Resource is not mapped! "
           "Resource should be mapped before accessing its memory.");
     }
 
@@ -1717,6 +1734,41 @@ class bindless_image_wrapper_base {
     return _img;
   }
 
+  /// Set mipmap sample filtering mode for bindless image handle
+  /// \param [in] filtering_mode The mipmap filtering mode
+  inline void set_mip_filtering_mode(sycl::filtering_mode mode) {
+    auto sampling_info = get_sampling_info(_img);
+    sampling_info.set_mipmap_filtering(mode);
+  }
+
+  /// Get mipmap sample filtering mode for bindless image handle
+  /// \return The mipmap filtering mode
+  inline sycl::filtering_mode get_mip_filtering_mode(void) {
+    auto sampling_info = get_sampling_info(_img);
+    return sampling_info.get_mipmap_filtering();
+  }
+
+  /// Get mipmap sample filtering mode for bindless image handle
+  /// \return The mipmap filtering mode
+  inline void get_mip_level_clamp(float *min_level_clamp,
+                                  float *max_level_clamp) {
+    auto sampling_info = get_sampling_info(_img);
+    *min_level_clamp = sampling_info.get_min_mipmap_level_clamp();
+    *max_level_clamp = sampling_info.get_max_mipmap_level_clamp();
+  }
+
+  /// Get mipmap memory wrapper attached the bindless image
+  /// \return The mipmap memory wrapper
+  inline image_mem_wrapper *get_attached_mipmap_data(void) {
+    auto mem = detail::get_img_mem_map(_img);
+
+    if (mem->get_image_type() !=
+        sycl::ext::oneapi::experimental::image_type::mipmap)
+      throw std::runtime_error("Bindless image data is not of mipmap type");
+
+    return mem;
+  }
+
 private:
   image_channel _channel;
   sycl::addressing_mode _addressing_mode = sycl::addressing_mode::clamp_to_edge;

clang/test/dpct/texture/texture_object_bindless_image.cu

@@ -654,3 +654,62 @@ int main() {
   cudaFreeMipmappedArray(pMipMapArr);
   return 0;
 }
+
+// CHECK: void set_3D_descriptor(sycl::ext::oneapi::experimental::image_descriptor &desc) {
+void set_3D_descriptor(CUDA_ARRAY3D_DESCRIPTOR &desc) {
+  desc.Width = 1;
+  desc.Depth = 2;
+  desc.Height = 1;
+  // CHECK: desc.channel_type = sycl::image_channel_type::unsigned_int16;
+  desc.Format = CU_AD_FORMAT_SIGNED_INT16;
+  desc.NumChannels = 2;
+}
+
+void test_mipmap_driver_api() {
+  // CHECK: sycl::ext::oneapi::experimental::image_descriptor desc;
+  CUDA_ARRAY3D_DESCRIPTOR desc;
+  unsigned int numMipmapLevels = 2;
+  set_3D_descriptor(desc);
+
+  // CHECK: dpct::experimental::image_mem_wrapper_ptr mmArray;
+  CUmipmappedArray mmArray;
+
+  // CHECK: mmArray = new dpct::experimental::image_mem_wrapper(desc, numMipmapLevels);
+  cuMipmappedArrayCreate(&mmArray, &desc, numMipmapLevels);
+
+  // CHECK: dpct::experimental::image_mem_wrapper_ptr *pArray;
+  CUmipmappedArray *pArray;
+  // CHECK: *pArray = new dpct::experimental::image_mem_wrapper(desc, numMipmapLevels);
+  cuMipmappedArrayCreate(pArray, &desc, numMipmapLevels);
+
+  CUarray level_arr;
+  // CHECK: level_arr = mmArray->get_mip_level(1);
+  cuMipmappedArrayGetLevel(&level_arr, mmArray, 1);
+
+  CUtexref texRef;
+  // CHECK: texRef->attach(mmArray);
+  cuTexRefSetMipmappedArray(texRef, mmArray, 0);
+
+  // sycl::filter_mode fm = sycl::filtering_mode::nearest;
+  CUfilter_mode fm = CU_TR_FILTER_MODE_POINT;
+
+  // CHECK: texRef->set_mip_filtering_mode(fm);
+  cuTexRefSetMipmapFilterMode(texRef, fm);
+
+  // CHECK: fm = texRef->get_mip_filtering_mode();
+  cuTexRefGetMipmapFilterMode(&fm, texRef);
+
+  float min_clamp, max_clamp;
+  // CHECK: texRef->get_mip_level_clamp(&min_clamp, &max_clamp);
+  cuTexRefGetMipmapLevelClamp(&min_clamp, &max_clamp, texRef);
+
+  CUmipmappedArray anotherArray;
+  // CHECK: anotherArray = texRef->get_attached_mipmap_data();
+  cuTexRefGetMipmappedArray(&anotherArray, texRef);
+
+  // CHECK: delete mmArray;
+  cuMipmappedArrayDestroy(mmArray);
+
+  // CHECK: delete (*pArray);
+  cuMipmappedArrayDestroy(*pArray);
+}

clang/test/dpct/texture_object.cu

@@ -437,4 +437,55 @@ void mipmap() {
   // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaTextureDesc::maxMipmapLevelClamp is not supported.
   // CHECK-NEXT: */
   texDesc.maxMipmapLevelClamp = 1;
+
+  CUDA_ARRAY3D_DESCRIPTOR texdesc;
+  unsigned int numMipmapLevels = 2;
+
+  CUmipmappedArray *pArray;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuMipmappedArrayCreate was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuMipmappedArrayCreate(pArray, &texdesc, numMipmapLevels);
+
+  CUarray level_arr;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuMipmappedArrayGetLevel was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuMipmappedArrayGetLevel(&level_arr, *pArray, 1);
+
+  CUtexref texRef;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuTexRefSetMipmappedArray was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuTexRefSetMipmappedArray(texRef, *pArray, 0);
+
+  // sycl::filter_mode fm = sycl::filtering_mode::nearest;
+  CUfilter_mode fm = CU_TR_FILTER_MODE_POINT;
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuTexRefSetMipmapFilterMode was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuTexRefSetMipmapFilterMode(texRef, fm);
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuTexRefGetMipmapFilterMode was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuTexRefGetMipmapFilterMode(&fm, texRef);
+
+  float min_clamp, max_clamp;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuTexRefGetMipmapLevelClamp was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuTexRefGetMipmapLevelClamp(&min_clamp, &max_clamp, texRef);
+
+  CUmipmappedArray anotherArray;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuTexRefGetMipmappedArray was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuTexRefGetMipmappedArray(&anotherArray, texRef);
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to cuMipmappedArrayDestroy was removed because SYCL currently does not support mipmap image type. You can migrate the code with bindless images by specifying --use-experimental-features=bindless_images.
+  // CHECK-NEXT: */
+  cuMipmappedArrayDestroy(*pArray);
 }