diff --git a/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc b/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc index f70efb70e9de..723f289cc0d8 100644 --- a/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc +++ b/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc @@ -107,10 +107,10 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY( // External Resource APIs CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, - CALL_FACTORY_ENTRY("cudaImportExternalMemory", + ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaImportExternalMemory", CALL(MapNames::getDpctNamespace() + "experimental::import_external_memory", - ARG(0), ARG(1))), + ARG(0), ARG(1)))), UNSUPPORT_FACTORY_ENTRY("cudaImportExternalMemory", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cudaImportExternalMemory"), ARG("--use-experimental-features=bindless_images"))) diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 41e7e28ddf66..e8dbf21b9829 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -122,7 +122,15 @@ class external_mem_img_desc { /// size setter /// \param [in] img_dims The dimensions of the imported image resource - void set_size(sycl::range<3> img_dims) { size = img_dims; } + void set_size(sycl::range<3> img_dims) { + size = img_dims; + + for (int i = 0; i < 3; i++) { + if (size[i] == 1) { + size[i] = 0; + } + } + } /// image_channel setter /// \param [in] img_ch The channel info of the imported image resource @@ -130,7 +138,13 @@ class external_mem_img_desc { /// num_levels setter /// \param [in] numLevels The no. of levels in the imported image resource - void set_num_levels(unsigned int numLevels) { num_levels = numLevels; } + void set_num_levels(unsigned int num_levels) { + this->num_levels = num_levels; + + if (num_levels > 1) { + set_image_type(sycl::ext::oneapi::experimental::image_type::mipmap); + } + } /// type setter /// \param [in] img_type The type of the imported image resource @@ -303,6 +317,11 @@ class image_mem_wrapper { /// Get the image mip level of the bindless image memory. /// \returns The image mip level of the bindless image memory. image_mem_wrapper *get_mip_level(unsigned int level) { + if (_desc.type == sycl::ext::oneapi::experimental::image_type::standard) { + assert(level == 0); + return this; + } + assert(_desc.type == sycl::ext::oneapi::experimental::image_type::mipmap); return _sub_wrappers + level; }