Skip to content

Commit ef25232

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

File tree

6 files changed

+63
-8
lines changed

6 files changed

+63
-8
lines changed

‎clang/lib/DPCT/APINames.inc‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -379,9 +379,9 @@ ENTRY(cudaGetTextureObjectResourceViewDesc, cudaGetTextureObjectResourceViewDesc
379379
ENTRY(cudaGetTextureObjectTextureDesc, cudaGetTextureObjectTextureDesc, true, NO_FLAG, P4, "Successful")
380380

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

386386
// Version Management
387387
ENTRY(cudaDriverGetVersion, cudaDriverGetVersion, true, NO_FLAG, P0, "DPCT1043")

‎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
@@ -12851,10 +12851,12 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
1285112851
)
1285212852
.bind("tex"),
1285312853
this);
12854-
MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
12855-
"cudaTextureObject_t", "CUtexObject"))))))
12856-
.bind("texObj"),
12857-
this);
12854+
MF.addMatcher(
12855+
typeLoc(
12856+
loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
12857+
"cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
12858+
.bind("texObj"),
12859+
this);
1285812860
MF.addMatcher(
1285912861
memberExpr(hasObjectExpression(hasType(
1286012862
type(hasUnqualifiedDesugaredType(recordType(hasDeclaration(
@@ -12914,6 +12916,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
1291412916
"cudaGetTextureObjectResourceDesc",
1291512917
"cudaGetTextureObjectTextureDesc",
1291612918
"cudaGetTextureObjectResourceViewDesc",
12919+
"cudaCreateSurfaceObject",
12920+
"cudaDestroySurfaceObject",
12921+
"cudaGetSurfaceObjectResourceDesc",
1291712922
"cuArray3DCreate_v2",
1291812923
"cuArrayCreate_v2",
1291912924
"cuArrayDestroy",

‎clang/lib/DPCT/MapNames.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -552,6 +552,13 @@ void MapNames::setExplicitNamespaceMap(
552552
"ext::oneapi::experimental::sampled_image_handle"
553553
: getDpctNamespace() + "image_wrapper_base_p",
554554
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)},
555562
{"textureReference",
556563
std::make_shared<TypeNameRule>(getDpctNamespace() + "image_wrapper_base",
557564
HelperFeatureEnum::device_ext)},
@@ -866,6 +873,7 @@ void MapNames::setExplicitNamespaceMap(
866873
"cudaTextureDesc",
867874
"cudaResourceDesc",
868875
"cudaTextureObject_t",
876+
"cudaSurfaceObject_t",
869877
"textureReference",
870878
"cudaTextureAddressMode",
871879
"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)