diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index a08b1fedac39..2b84f371c87f 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -554,6 +554,13 @@ void MapNames::setExplicitNamespaceMap( "ext::oneapi::experimental::sampled_image_handle" : getDpctNamespace() + "image_wrapper_base_p", HelperFeatureEnum::device_ext)}, + {"cudaSurfaceObject_t", + std::make_shared( + DpctGlobalInfo::useExtBindlessImages() + ? getClNamespace() + + "ext::oneapi::experimental::sampled_image_handle" + : getDpctNamespace() + "image_wrapper_base_p", + HelperFeatureEnum::device_ext)}, {"textureReference", std::make_shared(getDpctNamespace() + "image_wrapper_base", HelperFeatureEnum::device_ext)}, @@ -868,6 +875,7 @@ void MapNames::setExplicitNamespaceMap( "cudaTextureDesc", "cudaResourceDesc", "cudaTextureObject_t", + "cudaSurfaceObject_t", "textureReference", "cudaTextureAddressMode", "cudaTextureFilterMode", diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index ec2139dee9b4..d0ccccd864ee 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/clang/lib/DPCT/RulesLang/APINamesTexture.inc @@ -64,6 +64,13 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY( "experimental::destroy_bindless_image", ARG(0), QUEUESTR)), DELETER_FACTORY_ENTRY("cuTexObjectDestroy", ARG(0)))) +ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY( + UseExtBindlessImages, + CALL_FACTORY_ENTRY("cudaDestroySurfaceObject", + CALL(MapNames::getDpctNamespace() + + "experimental::destroy_bindless_image", + ARG(0), QUEUESTR)), + DELETER_FACTORY_ENTRY("cudaDestroySurfaceObject", ARG(0)))) CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, @@ -104,6 +111,15 @@ CONDITIONAL_FACTORY_ENTRY( "cuTexObjectGetTextureDesc", DEREF(0), MEMBER_CALL(ARG(1), true, "get_sampling_info"))))) +CONDITIONAL_FACTORY_ENTRY( + UseExtBindlessImages, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cudaGetSurfaceObjectResourceDesc", DEREF(0), + CALL(MapNames::getDpctNamespace() + "experimental::get_data", ARG(1)))), + FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cudaGetSurfaceObjectResourceDesc", DEREF(0), + MEMBER_CALL(ARG(1), true, "get_data"))))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ENTRY_TEXTURE("tex1D", 0x01, 1)) @@ -184,6 +200,16 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, "create_image_wrapper", DEREF(1), DEREF(2))))) +FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cudaCreateSurfaceObject", DEREF(0), + CALL(DpctGlobalInfo::useExtBindlessImages() + ? MapNames::getDpctNamespace() + + "experimental::create_bindless_image" + : MapNames::getDpctNamespace() + + "create_image_wrapper", + DEREF(1))))) + ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc", Diagnostics::API_NOT_MIGRATED) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 525af6316864..68cb87449ff4 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -9698,6 +9698,12 @@ void TextureRule::registerMatcher(MatchFinder &MF) { ) .bind("tex"), this); + MF.addMatcher( + typeLoc( + loc(qualType(hasDeclaration(typedefDecl(hasAnyName( + "cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject")))))) + .bind("texObj"), + this); MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName( "cudaTextureObject_t", "CUtexObject")))))) .bind("texObj"), @@ -9761,6 +9767,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) { "cudaGetTextureObjectResourceDesc", "cudaGetTextureObjectTextureDesc", "cudaGetTextureObjectResourceViewDesc", + "cudaCreateSurfaceObject", + "cudaDestroySurfaceObject", + "cudaGetSurfaceObjectResourceDesc", "cuArray3DCreate_v2", "cuArrayCreate_v2", "cuArrayDestroy", diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index f2fa3be225ca..93cccc99fbde 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/clang/lib/DPCT/SrcAPI/APINames.inc @@ -379,9 +379,9 @@ ENTRY(cudaGetTextureObjectResourceViewDesc, cudaGetTextureObjectResourceViewDesc ENTRY(cudaGetTextureObjectTextureDesc, cudaGetTextureObjectTextureDesc, true, NO_FLAG, P4, "Successful") // low level texture surface management functions of runtime API -ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, false, NO_FLAG, P0, "comment") -ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, false, NO_FLAG, P0, "comment") -ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, false, NO_FLAG, P4, "comment") +ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, true, NO_FLAG, P0, "Successful") +ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, true, NO_FLAG, P0, "Successful") +ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, true, NO_FLAG, P4, "Successful") // Version Management ENTRY(cudaDriverGetVersion, cudaDriverGetVersion, true, NO_FLAG, P0, "DPCT1043") diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 0df49d437e6a..0bbd4baa6691 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -844,7 +844,7 @@ inline void unmap_resources(int count, external_mem_wrapper **handles, /// \param [in] q The queue where the image creation be executed. /// \returns The sampled image handle of created bindless image. static inline sycl::ext::oneapi::experimental::sampled_image_handle -create_bindless_image(image_data data, sampling_info info, +create_bindless_image(image_data data, sampling_info info = {}, sycl::queue q = get_default_queue()) { auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler( info.get_addressing_mode(), info.get_coordinate_normalization_mode(), diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu new file mode 100644 index 000000000000..6a34429c63c7 --- /dev/null +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -0,0 +1,16 @@ +// RUN: dpct --format-range=none --use-experimental-features=bindless_images -out-root %T/texture/surface_object_bindless_image %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++14 +// RUN: FileCheck --input-file %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.cpp -o %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.o %} + +int main() { + // CHECK: sycl::ext::oneapi::experimental::sampled_image_handle surf; + cudaSurfaceObject_t surf; + // CHECK: dpct::image_data resDesc; + cudaResourceDesc resDesc; + // CHECK: surf = dpct::experimental::create_bindless_image(resDesc); + cudaCreateSurfaceObject(&surf, &resDesc); + // CHECK: dpct::experimental::destroy_bindless_image(surf, dpct::get_in_order_queue()); + cudaDestroySurfaceObject(surf); + // CHECK: resDesc = dpct::experimental::get_data(surf); + cudaGetSurfaceObjectResourceDesc(&resDesc, surf); +}