Skip to content

Commit e1b4fa3

Browse files
tangjj11tangjj11
committed
[SYCLomatic] Support migration of 3 runtime Surface APIs.
Signed-off-by: Tang, Jiajun [email protected]
1 parent 5b9bc9d commit e1b4fa3

File tree

6 files changed

+72
-13
lines changed

6 files changed

+72
-13
lines changed

‎clang/lib/DPCT/APINames.inc‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
66
//
77
//===----------------------------------------------------------------------===//
8+
// clang-format off
89
/****************************************************************************
910
*
1011
* ENTRY(Interface APIName, APINAME, ProcessedOrNot, Flag, Priority, MigrationDesc)
@@ -377,9 +378,9 @@ ENTRY(cudaGetTextureObjectResourceViewDesc, cudaGetTextureObjectResourceViewDesc
377378
ENTRY(cudaGetTextureObjectTextureDesc, cudaGetTextureObjectTextureDesc, true, NO_FLAG, P4, "Successful")
378379

379380
// low level texture surface management functions of runtime API
380-
ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, false, NO_FLAG, P0, "comment")
381-
ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, false, NO_FLAG, P0, "comment")
382-
ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, false, NO_FLAG, P4, "comment")
381+
ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, true, NO_FLAG, P0, "Successful")
382+
ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, true, NO_FLAG, P0, "Successful")
383+
ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, true, NO_FLAG, P4, "Successful")
383384

384385
// Version Management
385386
ENTRY(cudaDriverGetVersion, cudaDriverGetVersion, true, NO_FLAG, P0, "DPCT1043")
@@ -2307,3 +2308,4 @@ ENTRY(__assert_fail, __assert_fail, true, NO_FLAG, P4, "Successful")
23072308
ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful")
23082309

23092310
ENTRY(cuGetExportTable, cuGetExportTable, true, NO_FLAG, P7, "Partial")
2311+
// clang-format on

‎clang/lib/DPCT/APINamesTexture.inc‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,6 +64,13 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
6464
"experimental::destroy_bindless_image",
6565
ARG(0), QUEUESTR)),
6666
DELETER_FACTORY_ENTRY("cuTexObjectDestroy", ARG(0))))
67+
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
68+
UseExtBindlessImages,
69+
CALL_FACTORY_ENTRY("cudaDestroySurfaceObject",
70+
CALL(MapNames::getDpctNamespace() +
71+
"experimental::destroy_bindless_image",
72+
ARG(0), QUEUESTR)),
73+
DELETER_FACTORY_ENTRY("cudaDestroySurfaceObject", ARG(0))))
6774

6875
CONDITIONAL_FACTORY_ENTRY(
6976
UseExtBindlessImages,
@@ -104,6 +111,15 @@ CONDITIONAL_FACTORY_ENTRY(
104111
"cuTexObjectGetTextureDesc", DEREF(0),
105112
MEMBER_CALL(ARG(1), true,
106113
"get_sampling_info")))))
114+
CONDITIONAL_FACTORY_ENTRY(
115+
UseExtBindlessImages,
116+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
117+
"cudaGetSurfaceObjectResourceDesc", DEREF(0),
118+
CALL(MapNames::getDpctNamespace() + "experimental::get_data", ARG(1)))),
119+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
120+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
121+
"cudaGetSurfaceObjectResourceDesc", DEREF(0),
122+
MEMBER_CALL(ARG(1), true, "get_data")))))
107123

108124
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
109125
ENTRY_TEXTURE("tex1D", 0x01, 1))
@@ -184,6 +200,16 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
184200
"create_image_wrapper",
185201
DEREF(1), DEREF(2)))))
186202

203+
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
204+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
205+
"cudaCreateSurfaceObject", DEREF(0),
206+
CALL(DpctGlobalInfo::useExtBindlessImages()
207+
? MapNames::getDpctNamespace() +
208+
"experimental::create_bindless_image"
209+
: MapNames::getDpctNamespace() +
210+
"create_image_wrapper",
211+
DEREF(1)))))
212+
187213
ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
188214
Diagnostics::API_NOT_MIGRATED)
189215

‎clang/lib/DPCT/ASTTraversal.cpp‎

Lines changed: 9 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12840,10 +12840,12 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
1284012840
)
1284112841
.bind("tex"),
1284212842
this);
12843-
MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
12844-
"cudaTextureObject_t", "CUtexObject"))))))
12845-
.bind("texObj"),
12846-
this);
12843+
MF.addMatcher(
12844+
typeLoc(
12845+
loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
12846+
"cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
12847+
.bind("texObj"),
12848+
this);
1284712849
MF.addMatcher(
1284812850
memberExpr(hasObjectExpression(hasType(
1284912851
type(hasUnqualifiedDesugaredType(recordType(hasDeclaration(
@@ -12903,6 +12905,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
1290312905
"cudaGetTextureObjectResourceDesc",
1290412906
"cudaGetTextureObjectTextureDesc",
1290512907
"cudaGetTextureObjectResourceViewDesc",
12908+
"cudaCreateSurfaceObject",
12909+
"cudaDestroySurfaceObject",
12910+
"cudaGetSurfaceObjectResourceDesc",
1290612911
"cuArray3DCreate_v2",
1290712912
"cuArrayCreate_v2",
1290812913
"cuArrayDestroy",

‎clang/lib/DPCT/MapNames.cpp‎

Lines changed: 15 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -449,14 +449,16 @@ void MapNames::setExplicitNamespaceMap(
449449
getLibraryHelperNamespace() +
450450
"sparse::optimize_info>",
451451
HelperFeatureEnum::device_ext)},
452-
{"thrust::device_ptr",
453-
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_pointer",
454-
HelperFeatureEnum::device_ext)},
452+
{"thrust::device_ptr", std::make_shared<TypeNameRule>(
453+
getLibraryHelperNamespace() + "device_pointer",
454+
HelperFeatureEnum::device_ext)},
455455
{"thrust::device_reference",
456-
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_reference",
456+
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() +
457+
"device_reference",
457458
HelperFeatureEnum::device_ext)},
458459
{"thrust::device_vector",
459-
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() + "device_vector",
460+
std::make_shared<TypeNameRule>(getLibraryHelperNamespace() +
461+
"device_vector",
460462
HelperFeatureEnum::device_ext)},
461463
{"thrust::device_malloc_allocator",
462464
std::make_shared<TypeNameRule>(getDpctNamespace() +
@@ -550,6 +552,13 @@ void MapNames::setExplicitNamespaceMap(
550552
"ext::oneapi::experimental::sampled_image_handle"
551553
: getDpctNamespace() + "image_wrapper_base_p",
552554
HelperFeatureEnum::device_ext)},
555+
{"cudaSurfaceObject_t",
556+
std::make_shared<TypeNameRule>(
557+
DpctGlobalInfo::useExtBindlessImages()
558+
? getClNamespace() +
559+
"ext::oneapi::experimental::sampled_image_handle"
560+
: getDpctNamespace() + "image_wrapper_base_p",
561+
HelperFeatureEnum::device_ext)},
553562
{"textureReference",
554563
std::make_shared<TypeNameRule>(getDpctNamespace() + "image_wrapper_base",
555564
HelperFeatureEnum::device_ext)},
@@ -846,6 +855,7 @@ void MapNames::setExplicitNamespaceMap(
846855
"cudaTextureDesc",
847856
"cudaResourceDesc",
848857
"cudaTextureObject_t",
858+
"cudaSurfaceObject_t",
849859
"textureReference",
850860
"cudaTextureAddressMode",
851861
"cudaTextureFilterMode",

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -844,7 +844,7 @@ inline void unmap_resources(int count, external_mem_wrapper **handles,
844844
/// \param [in] q The queue where the image creation be executed.
845845
/// \returns The sampled image handle of created bindless image.
846846
static inline sycl::ext::oneapi::experimental::sampled_image_handle
847-
create_bindless_image(image_data data, sampling_info info,
847+
create_bindless_image(image_data data, sampling_info info = {},
848848
sycl::queue q = get_default_queue()) {
849849
auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
850850
info.get_addressing_mode(), info.get_coordinate_normalization_mode(),

‎clang/test/dpct/texture/surface_object_bindless_image.cu‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
1+
// 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
2+
// RUN: FileCheck --input-file %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.cpp --match-full-lines %s
3+
// 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 %}
4+
5+
int main() {
6+
// CHECK: sycl::ext::oneapi::experimental::sampled_image_handle surf;
7+
cudaSurfaceObject_t surf;
8+
// CHECK: dpct::image_data resDesc;
9+
cudaResourceDesc resDesc;
10+
// CHECK: surf = dpct::experimental::create_bindless_image(resDesc);
11+
cudaCreateSurfaceObject(&surf, &resDesc);
12+
// CHECK: dpct::experimental::destroy_bindless_image(surf, dpct::get_in_order_queue());
13+
cudaDestroySurfaceObject(surf);
14+
// CHECK: resDesc = dpct::experimental::get_data(surf);
15+
cudaGetSurfaceObjectResourceDesc(&resDesc, surf);
16+
}

0 commit comments

Comments
 (0)