From c0bc12774476ea28c1b6bc211a70c5bb2c16411d Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 9 Jul 2025 14:07:09 +0800 Subject: [PATCH 1/2] [SYCLomatic] Add 11 cuda driver API query Mappings. Signed-off-by: Chen, Sheng S --- .../cuMemGetAllocationPropertiesFromHandle.cu | 9 ++++ .../DPCT/Driver/cuMipmappedArrayDestroy.cu | 7 +++ .../DPCT/Driver/cuMipmappedArrayGetLevel.cu | 8 ++++ .../DPCT/Driver/cuPointerGetAttribute.cu | 9 ++++ clang/examples/DPCT/Driver/cuStreamGetCtx.cu | 6 +++ clang/examples/DPCT/Driver/cuTexRefCreate.cu | 8 ++++ clang/examples/DPCT/Driver/cuTexRefDestroy.cu | 8 ++++ .../examples/DPCT/Driver/cuTexRefGetArray.cu | 9 ++++ .../examples/DPCT/Driver/cuTexRefGetFlags.cu | 1 + .../Driver/cuTexRefGetMipmapFilterMode.cu | 9 ++++ .../Driver/cuTexRefGetMipmapLevelClamp.cu | 10 +++++ clang/lib/DPCT/DPCT.cpp | 3 ++ clang/lib/DPCT/RulesLang/APINamesTexture.inc | 11 +++++ clang/lib/DPCT/RulesLang/RulesLangTexture.cpp | 1 + .../query_api_mapping/Driver/test-after9.2.cu | 27 +++++++++++ .../dpct/query_api_mapping/Driver/test.cu | 45 +++++++++++++++++++ clang/test/dpct/query_api_mapping/test_all.cu | 10 +++++ .../texture_reference_bindless_image.cu | 2 + 18 files changed, 183 insertions(+) create mode 100644 clang/examples/DPCT/Driver/cuMemGetAllocationPropertiesFromHandle.cu create mode 100644 clang/examples/DPCT/Driver/cuMipmappedArrayDestroy.cu create mode 100644 clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu create mode 100644 clang/examples/DPCT/Driver/cuPointerGetAttribute.cu create mode 100644 clang/examples/DPCT/Driver/cuStreamGetCtx.cu create mode 100644 clang/examples/DPCT/Driver/cuTexRefCreate.cu create mode 100644 clang/examples/DPCT/Driver/cuTexRefDestroy.cu create mode 100644 clang/examples/DPCT/Driver/cuTexRefGetArray.cu create mode 100644 clang/examples/DPCT/Driver/cuTexRefGetMipmapFilterMode.cu create mode 100644 clang/examples/DPCT/Driver/cuTexRefGetMipmapLevelClamp.cu diff --git a/clang/examples/DPCT/Driver/cuMemGetAllocationPropertiesFromHandle.cu b/clang/examples/DPCT/Driver/cuMemGetAllocationPropertiesFromHandle.cu new file mode 100644 index 000000000000..1f88686bc19f --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMemGetAllocationPropertiesFromHandle.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=virtual_mem +#include +void test() { + CUmemAllocationProp prop = {}; + CUmemGenericAllocationHandle allocHandle; + // Start + cuMemGetAllocationPropertiesFromHandle(&prop/*CUmemAllocationProp **/, allocHandle/*CUmemGenericAllocationHandle*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMipmappedArrayDestroy.cu b/clang/examples/DPCT/Driver/cuMipmappedArrayDestroy.cu new file mode 100644 index 000000000000..13553933518f --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMipmappedArrayDestroy.cu @@ -0,0 +1,7 @@ +// Option: --use-experimental-features=bindless_images +void test() { + CUmipmappedArray mmArray; + // Start + cuMipmappedArrayDestroy(mmArray/*CUmipmappedArray*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu b/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu new file mode 100644 index 000000000000..cf009c0ac808 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images +void test() { + CUmipmappedArray mmArray; + CUarray level_arr; + // Start + cuMipmappedArrayGetLevel(&level_arr/*CUarray*/, mmArray/*CUmipmappedArray*/, 1/*level*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuPointerGetAttribute.cu b/clang/examples/DPCT/Driver/cuPointerGetAttribute.cu new file mode 100644 index 000000000000..aedfbc709550 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuPointerGetAttribute.cu @@ -0,0 +1,9 @@ +// Option: --usm-level=none +#include +void test() { + void *base_ptr; + void *ptr; + // Start + cuPointerGetAttribute(base_ptr/*void **/, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR/*CUpointer_attribute*/, (CUdeviceptr)ptr/*CUdeviceptr*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuStreamGetCtx.cu b/clang/examples/DPCT/Driver/cuStreamGetCtx.cu new file mode 100644 index 000000000000..31a296f8c264 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuStreamGetCtx.cu @@ -0,0 +1,6 @@ +#include +void test(CUstream stream, CUcontext& context) { + // Start + cuStreamGetCtx(stream/*CUstream*/, &context/*CUcontext **/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuTexRefCreate.cu b/clang/examples/DPCT/Driver/cuTexRefCreate.cu new file mode 100644 index 000000000000..67ef151b4a6f --- /dev/null +++ b/clang/examples/DPCT/Driver/cuTexRefCreate.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images +#include +void test() { + CUtexref r; + // Start + cuTexRefCreate(&r/*CUtexref **/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuTexRefDestroy.cu b/clang/examples/DPCT/Driver/cuTexRefDestroy.cu new file mode 100644 index 000000000000..29e35ed6fb10 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuTexRefDestroy.cu @@ -0,0 +1,8 @@ +// Option: --use-experimental-features=bindless_images +#include +void test() { + CUtexref r; + // Start + cuTexRefDestroy(r/*CUtexref*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuTexRefGetArray.cu b/clang/examples/DPCT/Driver/cuTexRefGetArray.cu new file mode 100644 index 000000000000..02d4079dbee0 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuTexRefGetArray.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=bindless_images +#include +void test() { + CUtexref r; + CUarray a; + // Start + cuTexRefGetArray(&a/*CUarray **/, r/*CUtexref*/); + // End +} \ No newline at end of file diff --git a/clang/examples/DPCT/Driver/cuTexRefGetFlags.cu b/clang/examples/DPCT/Driver/cuTexRefGetFlags.cu index b8bf4198333a..ebca8ba1f7a0 100644 --- a/clang/examples/DPCT/Driver/cuTexRefGetFlags.cu +++ b/clang/examples/DPCT/Driver/cuTexRefGetFlags.cu @@ -1,3 +1,4 @@ +#include void test(unsigned int *pu) { // Start CUtexref t; diff --git a/clang/examples/DPCT/Driver/cuTexRefGetMipmapFilterMode.cu b/clang/examples/DPCT/Driver/cuTexRefGetMipmapFilterMode.cu new file mode 100644 index 000000000000..07ffc89bc88d --- /dev/null +++ b/clang/examples/DPCT/Driver/cuTexRefGetMipmapFilterMode.cu @@ -0,0 +1,9 @@ +// Option: --use-experimental-features=bindless_images +#include +void test() { + CUfilter_mode fm = CU_TR_FILTER_MODE_POINT; + CUtexref texRef; + // Start + cuTexRefGetMipmapFilterMode(&fm /*CUfilter_mode **/, texRef /*CUtexref*/); + // End +} diff --git a/clang/examples/DPCT/Driver/cuTexRefGetMipmapLevelClamp.cu b/clang/examples/DPCT/Driver/cuTexRefGetMipmapLevelClamp.cu new file mode 100644 index 000000000000..fe2b37831091 --- /dev/null +++ b/clang/examples/DPCT/Driver/cuTexRefGetMipmapLevelClamp.cu @@ -0,0 +1,10 @@ +// Option: --use-experimental-features=bindless_images +#include + +void test() { + CUtexref texRef; + float min_clamp, max_clamp; + // Start + cuTexRefGetMipmapLevelClamp(&min_clamp/*float **/, &max_clamp/*float **/, texRef/*CUtexref*/); + // End +} diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 38cb478308d1..ca70d9886928 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -1065,6 +1065,9 @@ int runDPCT(int argc, const char **argv) { NoDRYPattern.setValue(true); } else if (Option == "--enable-profiling") { EnablepProfiling.setValue(true); + } else if (Option.starts_with("--usm-level=")) { + if (Option.ends_with("none")) + USMLevel.setValue(UsmLevel::UL_None); } // Need add more option. } diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index 9e180b0512b9..ba8bb80b1d50 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/clang/lib/DPCT/RulesLang/APINamesTexture.inc @@ -467,6 +467,17 @@ CONDITIONAL_FACTORY_ENTRY( Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cuSurfRefGetArray"), ARG("--use-experimental-features=bindless_images"))) +CONDITIONAL_FACTORY_ENTRY( + UseExtBindlessImages, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cuTexRefGetArray", DEREF(0), + CALL(MapNames::getDpctNamespace() + "experimental::get_img_mem", + ARG(1)))), + UNSUPPORT_FACTORY_ENTRY("cuTexRefGetArray", + Diagnostics::TRY_EXPERIMENTAL_FEATURE, + ARG("cuTexRefGetArray"), + ARG("--use-experimental-features=bindless_images"))) + CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, diff --git a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp index 23a634b3595e..0ba5467e158d 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp @@ -589,6 +589,7 @@ void TextureRule::registerMatcher(MatchFinder &MF) { "cuTexRefDestroy", "cuSurfRefSetArray", "cuSurfRefGetArray", + "cuTexRefGetArray", "cuTexRefSetArray", "cuTexRefSetFormat", "cuTexRefSetAddressMode", diff --git a/clang/test/dpct/query_api_mapping/Driver/test-after9.2.cu b/clang/test/dpct/query_api_mapping/Driver/test-after9.2.cu index a7e4cc3212d6..c309fc69cc36 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test-after9.2.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test-after9.2.cu @@ -8,3 +8,30 @@ // CUDEVICEGETUUID-NEXT: cuDeviceGetUuid(pu /*CUuuid **/, d /*CUdevice*/); // CUDEVICEGETUUID-NEXT: Is migrated to: // CUDEVICEGETUUID-NEXT: *pu = dpct::get_device(d).get_device_info().get_uuid(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemAddressReserve | FileCheck %s -check-prefix=CUMEMADDRESSRESERVE +// CUMEMADDRESSRESERVE: CUDA API: +// CUMEMADDRESSRESERVE-NEXT: cuMemAddressReserve(ptr /*CUdeviceptr **/, size /*size_t*/, +// CUMEMADDRESSRESERVE-NEXT: alignment /*size_t*/, addr /*CUdeviceptr*/, +// CUMEMADDRESSRESERVE-NEXT: flags /*unsigned long long*/); +// CUMEMADDRESSRESERVE-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// CUMEMADDRESSRESERVE-NEXT: *ptr = (dpct::device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem((uintptr_t)addr, size, dpct::get_current_device().get_context()); + + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemGetAllocationPropertiesFromHandle | FileCheck %s -check-prefix=CUMEMGETALLOCATIONPROPERTIESFROMHANDLE +// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE: CUDA API: +// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT: cuMemGetAllocationPropertiesFromHandle(&prop/*CUmemAllocationProp **/, allocHandle/*CUmemGenericAllocationHandle*/); +// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem): +// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT: prop.location.id = dpct::get_device_id(allocHandle->get_device()); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttribute | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTE +// CUPOINTERGETATTRIBUTE: CUDA API: +// CUPOINTERGETATTRIBUTE-NEXT: cuPointerGetAttribute(base_ptr/*void **/, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR/*CUpointer_attribute*/, (CUdeviceptr)ptr/*CUdeviceptr*/); +// CUPOINTERGETATTRIBUTE-NEXT: Is migrated to (with the option --usm-level=none): +// CUPOINTERGETATTRIBUTE-NEXT: base_ptr = dpct::get_base_addr((dpct::device_ptr)ptr); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuStreamGetCtx | FileCheck %s -check-prefix=CUSTREAMGETCTX +// CUSTREAMGETCTX: CUDA API: +// CUSTREAMGETCTX-NEXT: cuStreamGetCtx(stream/*CUstream*/, &context/*CUcontext **/); +// CUSTREAMGETCTX-NEXT: Is migrated to: +// CUSTREAMGETCTX-NEXT: context = dpct::get_device_id(stream->get_device()); diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index ae871f9af460..01d73d512705 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -561,3 +561,48 @@ // CUDEVICECANACCESSPEER-NEXT: cuDeviceCanAccessPeer(pi /*int **/, d1 /*CUdevice*/, d2 /*CUdevice*/); // CUDEVICECANACCESSPEER-NEXT: Is migrated to: // CUDEVICECANACCESSPEER-NEXT: *pi = dpct::get_device(d1).ext_oneapi_can_access_peer(dpct::get_device(d2)); + + +/// Bindless Image + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMipmappedArrayDestroy | FileCheck %s -check-prefix=CUMIPMAPPEDARRAYDESTROY +// CUMIPMAPPEDARRAYDESTROY: CUDA API: +// CUMIPMAPPEDARRAYDESTROY-NEXT: cuMipmappedArrayDestroy(mmArray/*CUmipmappedArray*/); +// CUMIPMAPPEDARRAYDESTROY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUMIPMAPPEDARRAYDESTROY-NEXT: delete mmArray; + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMipmappedArrayGetLevel | FileCheck %s -check-prefix=CUMIPMAPPEDARRAYGETLEVEL +// CUMIPMAPPEDARRAYGETLEVEL: CUDA API: +// CUMIPMAPPEDARRAYGETLEVEL-NEXT: cuMipmappedArrayGetLevel(&level_arr/*CUarray*/, mmArray/*CUmipmappedArray*/, 1/*level*/); +// CUMIPMAPPEDARRAYGETLEVEL-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUMIPMAPPEDARRAYGETLEVEL-NEXT: level_arr = mmArray->get_mip_level(1); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefCreate | FileCheck %s -check-prefix=CUTEXREFCREATE +// CUTEXREFCREATE: CUDA API: +// CUTEXREFCREATE-NEXT: cuTexRefCreate(&r/*CUtexref **/); +// CUTEXREFCREATE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUTEXREFCREATE-NEXT: r = new dpct::experimental::bindless_image_wrapper_base(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefDestroy | FileCheck %s -check-prefix=CUTEXREFDESTROY +// CUTEXREFDESTROY: CUDA API: +// CUTEXREFDESTROY-NEXT: cuTexRefDestroy(r/*CUtexref*/); +// CUTEXREFDESTROY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUTEXREFDESTROY-NEXT: delete r; + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetArray | FileCheck %s -check-prefix=CUTEXREFGETARRAY +// CUTEXREFGETARRAY: CUDA API: +// CUTEXREFGETARRAY-NEXT: cuTexRefGetArray(phArray /*CUarray **/, hTexRef /*CUtexref*/); +// CUTEXREFGETARRAY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUTEXREFGETARRAY-NEXT: a = dpct::experimental::get_img_mem(r); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetMipmapFilterMode | FileCheck %s -check-prefix=CUTEXREFGETMIPMAPFILTERMODE +// CUTEXREFGETMIPMAPFILTERMODE: CUDA API: +// CUTEXREFGETMIPMAPFILTERMODE-NEXT: cuTexRefGetMipmapFilterMode(&fm /*CUfilter_mode **/, texRef /*CUtexref*/); +// CUTEXREFGETMIPMAPFILTERMODE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUTEXREFGETMIPMAPFILTERMODE-NEXT: fm = texRef->get_mip_filtering_mode(); + +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetMipmapLevelClamp | FileCheck %s -check-prefix=CUTEXREFGETMIPMAPLEVELCLAMP +// CUTEXREFGETMIPMAPLEVELCLAMP: CUDA API: +// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT: cuTexRefGetMipmapLevelClamp(&min_clamp/*float **/, &max_clamp/*float **/, texRef/*CUtexref*/); +// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images): +// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT: texRef->get_mip_level_clamp(&min_clamp, &max_clamp); \ No newline at end of file diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index 1eae79a9b5ea..07e850b7aa2f 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -757,6 +757,7 @@ // CHECK-NEXT: cuMemFree // CHECK-NEXT: cuMemFreeHost // CHECK-NEXT: cuMemGetAllocationGranularity +// CHECK-NEXT: cuMemGetAllocationPropertiesFromHandle // CHECK-NEXT: cuMemGetInfo // CHECK-NEXT: cuMemHostAlloc // CHECK-NEXT: cuMemHostGetDevicePointer @@ -804,6 +805,8 @@ // CHECK-NEXT: cuMemsetD32Async // CHECK-NEXT: cuMemsetD8 // CHECK-NEXT: cuMemsetD8Async +// CHECK-NEXT: cuMipmappedArrayDestroy +// CHECK-NEXT: cuMipmappedArrayGetLevel // CHECK-NEXT: cuModuleGetFunction // CHECK-NEXT: cuModuleGetTexRef // CHECK-NEXT: cuModuleLoad @@ -811,11 +814,13 @@ // CHECK-NEXT: cuModuleLoadDataEx // CHECK-NEXT: cuModuleUnload // CHECK-NEXT: cuOccupancyMaxActiveBlocksPerMultiprocessor +// CHECK-NEXT: cuPointerGetAttribute // CHECK-NEXT: cuPointerGetAttributes // CHECK-NEXT: cuStreamAddCallback // CHECK-NEXT: cuStreamAttachMemAsync // CHECK-NEXT: cuStreamCreate // CHECK-NEXT: cuStreamDestroy +// CHECK-NEXT: cuStreamGetCtx // CHECK-NEXT: cuStreamQuery // CHECK-NEXT: cuStreamSynchronize // CHECK-NEXT: cuStreamWaitEvent @@ -823,9 +828,14 @@ // CHECK-NEXT: cuTexObjectDestroy // CHECK-NEXT: cuTexObjectGetResourceDesc // CHECK-NEXT: cuTexObjectGetTextureDesc +// CHECK-NEXT: cuTexRefCreate +// CHECK-NEXT: cuTexRefDestroy // CHECK-NEXT: cuTexRefGetAddressMode +// CHECK-NEXT: cuTexRefGetArray // CHECK-NEXT: cuTexRefGetFilterMode // CHECK-NEXT: cuTexRefGetFlags +// CHECK-NEXT: cuTexRefGetMipmapFilterMode +// CHECK-NEXT: cuTexRefGetMipmapLevelClamp // CHECK-NEXT: cuTexRefSetAddress // CHECK-NEXT: cuTexRefSetAddress2D // CHECK-NEXT: cuTexRefSetAddressMode diff --git a/clang/test/dpct/texture/texture_reference_bindless_image.cu b/clang/test/dpct/texture/texture_reference_bindless_image.cu index 7ea9776d8f8a..54266f83801d 100644 --- a/clang/test/dpct/texture/texture_reference_bindless_image.cu +++ b/clang/test/dpct/texture/texture_reference_bindless_image.cu @@ -75,6 +75,8 @@ void driverTextureReferenceManagement() { // CHECK: r->set_channel_type(f); // CHECK-NEXT: r->set_channel_num(i); cuTexRefSetFormat(r, f, i); + + cuTexRefGetArray(&a, r); } void test_surf_ref() { // CHECK: dpct::experimental::image_mem_wrapper_ptr arr; From 1be0f09d7d83d1c3c91c12c93cedd8e8ed96b008 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 10 Jul 2025 09:15:41 +0800 Subject: [PATCH 2/2] up Signed-off-by: Chen, Sheng S --- clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu | 2 +- clang/test/dpct/texture/texture_reference_bindless_image.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu b/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu index cf009c0ac808..207f78e73393 100644 --- a/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu +++ b/clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu @@ -3,6 +3,6 @@ void test() { CUmipmappedArray mmArray; CUarray level_arr; // Start - cuMipmappedArrayGetLevel(&level_arr/*CUarray*/, mmArray/*CUmipmappedArray*/, 1/*level*/); + cuMipmappedArrayGetLevel(&level_arr/*CUarray **/, mmArray/*CUmipmappedArray*/, 1/*unsigned int*/); // End } \ No newline at end of file diff --git a/clang/test/dpct/texture/texture_reference_bindless_image.cu b/clang/test/dpct/texture/texture_reference_bindless_image.cu index 54266f83801d..350bee2235ef 100644 --- a/clang/test/dpct/texture/texture_reference_bindless_image.cu +++ b/clang/test/dpct/texture/texture_reference_bindless_image.cu @@ -75,7 +75,7 @@ void driverTextureReferenceManagement() { // CHECK: r->set_channel_type(f); // CHECK-NEXT: r->set_channel_num(i); cuTexRefSetFormat(r, f, i); - + // CHECK: a = dpct::experimental::get_img_mem(r); cuTexRefGetArray(&a, r); } void test_surf_ref() {