Skip to content

Commit 6e8daad

Browse files
ShengchenJShengchenJ
authored
[SYCLomatic] Add migration for 7 texture and surface APIs . (#2755)
Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>
1 parent 6957858 commit 6e8daad

File tree

9 files changed

+188
-36
lines changed

9 files changed

+188
-36
lines changed

‎clang/lib/DPCT/RuleInfra/MapNames.cpp‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -589,6 +589,13 @@ void MapNames::setExplicitNamespaceMap(
589589
"ext::oneapi::experimental::unsampled_image_handle"
590590
: getDpctNamespace() + "image_wrapper_base_p",
591591
HelperFeatureEnum::device_ext)},
592+
{"CUsurfref",
593+
std::make_shared<TypeNameRule>(
594+
DpctGlobalInfo::useExtBindlessImages()
595+
? getClNamespace() +
596+
"ext::oneapi::experimental::unsampled_image_handle"
597+
: getDpctNamespace() + "image_wrapper_base_p",
598+
HelperFeatureEnum::device_ext)},
592599
{"textureReference",
593600
std::make_shared<TypeNameRule>(getDpctNamespace() + "image_wrapper_base",
594601
HelperFeatureEnum::device_ext)},

‎clang/lib/DPCT/RulesLang/APINamesTexture.inc‎

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -295,6 +295,26 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
295295
ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
296296
Diagnostics::API_NOT_MIGRATED)
297297

298+
CONDITIONAL_FACTORY_ENTRY(
299+
UseExtBindlessImages,
300+
ASSIGNABLE_FACTORY(
301+
ASSIGN_FACTORY_ENTRY("cuArray3DGetDescriptor_v2", DEREF(0),
302+
MEMBER_CALL(ARG(1), true, "get_desc"))),
303+
UNSUPPORT_FACTORY_ENTRY("cuArray3DGetDescriptor_v2",
304+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
305+
ARG("cuArray3DGetDescriptor"),
306+
ARG("--use-experimental-features=bindless_images")))
307+
308+
CONDITIONAL_FACTORY_ENTRY(
309+
UseExtBindlessImages,
310+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cuArrayGetDescriptor_v2", DEREF(0),
311+
MEMBER_CALL(ARG(1), true,
312+
"get_desc"))),
313+
UNSUPPORT_FACTORY_ENTRY("cuArrayGetDescriptor_v2",
314+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
315+
ARG("cuArrayGetDescriptor"),
316+
ARG("--use-experimental-features=bindless_images")))
317+
298318
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
299319
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
300320
"cuArray3DCreate_v2", DEREF(0),
@@ -424,6 +444,41 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
424444
"cuTexRefSetAddress2D_v3", ARG(0), true, "attach",
425445
ARG(1), ARG(2), ARG(3))))
426446

447+
CONDITIONAL_FACTORY_ENTRY(
448+
UseExtBindlessImages,
449+
ASSIGNABLE_FACTORY(
450+
ASSIGN_FACTORY_ENTRY("cuTexRefCreate", DEREF(0),
451+
NEW(MapNames::getDpctNamespace() +
452+
"experimental::bindless_image_wrapper_base"))),
453+
UNSUPPORT_FACTORY_ENTRY("cuTexRefCreate",
454+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
455+
ARG("cuTexRefCreate"),
456+
ARG("--use-experimental-features=bindless_images")))
457+
458+
ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuTexRefDestroy", ARG(0)))
459+
460+
CONDITIONAL_FACTORY_ENTRY(
461+
UseExtBindlessImages,
462+
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
463+
"cuSurfRefGetArray", DEREF(0),
464+
CALL(MapNames::getDpctNamespace() + "experimental::get_img_mem",
465+
ARG(1)))),
466+
UNSUPPORT_FACTORY_ENTRY("cuSurfRefGetArray",
467+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
468+
ARG("cuSurfRefGetArray"),
469+
ARG("--use-experimental-features=bindless_images")))
470+
471+
CONDITIONAL_FACTORY_ENTRY(
472+
UseExtBindlessImages,
473+
CALL_FACTORY_ENTRY("cuSurfRefSetArray",
474+
CALL(MapNames::getDpctNamespace() +
475+
"experimental::set_img_mem",
476+
ARG(0), ARG(1))),
477+
UNSUPPORT_FACTORY_ENTRY("cuSurfRefSetArray",
478+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
479+
ARG("cuSurfRefSetArray"),
480+
ARG("--use-experimental-features=bindless_images")))
481+
427482
FEATURE_REQUEST_FACTORY(
428483
HelperFeatureEnum::device_ext,
429484
MULTI_STMTS_FACTORY_ENTRY(

‎clang/lib/DPCT/RulesLang/RulesLang.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -348,7 +348,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
348348
"cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
349349
"cudaExternalSemaphoreHandleType", "CUstreamCallback",
350350
"cudaHostFn_t", "__nv_half2", "__nv_half",
351-
"cudaGraphNodeType"))))))
351+
"cudaGraphNodeType", "CUsurfref"))))))
352352
.bind("cudaTypeDef"),
353353
this);
354354

‎clang/lib/DPCT/RulesLang/RulesLangTexture.cpp‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -575,6 +575,8 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
575575
"cuSurfObjectDestroy",
576576
"cuArray3DCreate_v2",
577577
"cuArrayCreate_v2",
578+
"cuArray3DGetDescriptor_v2",
579+
"cuArrayGetDescriptor_v2",
578580
"cuMipmappedArrayCreate",
579581
"cuMipmappedArrayDestroy",
580582
"cuMipmappedArrayGetLevel",
@@ -583,6 +585,10 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
583585
"cuTexObjectDestroy",
584586
"cuTexObjectGetTextureDesc",
585587
"cuTexObjectGetResourceDesc",
588+
"cuTexRefCreate",
589+
"cuTexRefDestroy",
590+
"cuSurfRefSetArray",
591+
"cuSurfRefGetArray",
586592
"cuTexRefSetArray",
587593
"cuTexRefSetFormat",
588594
"cuTexRefSetAddressMode",

‎clang/lib/DPCT/SrcAPI/APINames.inc‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1646,10 +1646,10 @@ ENTRY(cuLibraryUnload, cuLibraryUnload, false, NO_FLAG, P4, "comment")
16461646

16471647
// Memory Management
16481648
ENTRY(cuArray3DCreate, cuArray3DCreate_v2, true, NO_FLAG, P4, "Successful")
1649-
ENTRY(cuArray3DGetDescriptor, cuArray3DGetDescriptor_v2, false, NO_FLAG, P4, "comment")
1649+
ENTRY(cuArray3DGetDescriptor, cuArray3DGetDescriptor_v2, true, NO_FLAG, P4, "comment")
16501650
ENTRY(cuArrayCreate, cuArrayCreate_v2, true, NO_FLAG, P4, "Successful")
16511651
ENTRY(cuArrayDestroy, cuArrayDestroy, true, NO_FLAG, P4, "Successful")
1652-
ENTRY(cuArrayGetDescriptor, cuArrayGetDescriptor_v2, false, NO_FLAG, P4, "comment")
1652+
ENTRY(cuArrayGetDescriptor, cuArrayGetDescriptor_v2, true, NO_FLAG, P4, "comment")
16531653
ENTRY(cuArrayGetMemoryRequirements, cuArrayGetMemoryRequirements, false, NO_FLAG, P4, "comment")
16541654
ENTRY(cuArrayGetPlane, cuArrayGetPlane, false, NO_FLAG, P7, "comment")
16551655
ENTRY(cuArrayGetSparseProperties, cuArrayGetSparseProperties, false, NO_FLAG, P7, "comment")
@@ -1952,11 +1952,11 @@ ENTRY(cuOccupancyMaxPotentialBlockSizeWithFlags, cuOccupancyMaxPotentialBlockSiz
19521952
ENTRY(cuOccupancyMaxPotentialClusterSize, cuOccupancyMaxPotentialClusterSize, false, NO_FLAG, P4, "comment")
19531953

19541954
// Texture Reference Management(Deprecated)
1955-
ENTRY(cuTexRefCreate, cuTexRefCreate, false, NO_FLAG, P4, "comment")
1956-
ENTRY(cuTexRefDestroy, cuTexRefDestroy, false, NO_FLAG, P4, "comment")
1955+
ENTRY(cuTexRefCreate, cuTexRefCreate, true, NO_FLAG, P4, "comment")
1956+
ENTRY(cuTexRefDestroy, cuTexRefDestroy, true, NO_FLAG, P4, "comment")
19571957
ENTRY(cuTexRefGetAddress, cuTexRefGetAddress, false, NO_FLAG, P4, "comment")
19581958
ENTRY(cuTexRefGetAddressMode, cuTexRefGetAddressMode, true, NO_FLAG, P4, "Successful")
1959-
ENTRY(cuTexRefGetArray, cuTexRefGetArray, false, NO_FLAG, P4, "comment")
1959+
ENTRY(cuTexRefGetArray, cuTexRefGetArray, true, NO_FLAG, P4, "comment")
19601960
ENTRY(cuTexRefGetBorderColor, cuTexRefGetBorderColor, false, NO_FLAG, P4, "comment")
19611961
ENTRY(cuTexRefGetFilterMode, cuTexRefGetFilterMode, true, NO_FLAG, P4, "Successful")
19621962
ENTRY(cuTexRefGetFlags, cuTexRefGetFlags, true, NO_FLAG, P4, "Successful")
@@ -1981,8 +1981,8 @@ ENTRY(cuTexRefSetMipmapLevelClamp, cuTexRefSetMipmapLevelClamp, false, NO_FLAG,
19811981
ENTRY(cuTexRefSetMipmappedArray, cuTexRefSetMipmappedArray, true, NO_FLAG, P4, "Successful")
19821982

19831983
// Surface Reference Management(Deprecated)
1984-
ENTRY(cuSurfRefGetArray, cuSurfRefGetArray, false, NO_FLAG, P4, "comment")
1985-
ENTRY(cuSurfRefSetArray, cuSurfRefSetArray, false, NO_FLAG, P4, "comment")
1984+
ENTRY(cuSurfRefGetArray, cuSurfRefGetArray, true, NO_FLAG, P4, "comment")
1985+
ENTRY(cuSurfRefSetArray, cuSurfRefSetArray, true, NO_FLAG, P4, "comment")
19861986

19871987
// Texture Object Management
19881988
ENTRY(cuTexObjectCreate, cuTexObjectCreate, true, NO_FLAG, P4, "Successful")

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

Lines changed: 58 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -913,21 +913,7 @@ inline std::pair<image_data, sampling_info> &get_img_info_map(
913913
img_info_map;
914914
return img_info_map[handle];
915915
}
916-
inline image_mem_wrapper *&get_img_mem_map(
917-
const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
918-
static std::map<sycl::ext::oneapi::experimental::sampled_image_handle,
919-
image_mem_wrapper *, sampled_image_handle_compare>
920-
img_mem_map;
921-
return img_mem_map[handle];
922-
}
923916

924-
inline image_mem_wrapper *&get_img_mem_map(
925-
const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
926-
static std::map<sycl::ext::oneapi::experimental::unsampled_image_handle,
927-
image_mem_wrapper *, sampled_image_handle_compare>
928-
img_mem_map;
929-
return img_mem_map[handle];
930-
}
931917
static inline size_t
932918
get_ele_size(const sycl::ext::oneapi::experimental::image_descriptor &decs) {
933919
size_t channel_size;
@@ -1185,6 +1171,56 @@ inline bool check_duplicate_entries(int count, T **entries) {
11851171
}
11861172
} // namespace detail
11871173

1174+
/// Get image_mem_wrapper according to sampled image handle.
1175+
/// \param [in] handle The bindless image handle.
1176+
/// \returns The image_mem_wrapper of sampled image.
1177+
inline image_mem_wrapper *&get_img_mem(
1178+
const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
1179+
static std::map<sycl::ext::oneapi::experimental::sampled_image_handle,
1180+
image_mem_wrapper *, detail::sampled_image_handle_compare>
1181+
img_mem_map;
1182+
return img_mem_map[handle];
1183+
}
1184+
1185+
/// Get image_mem_wrapper according to unsampled image handle.
1186+
/// \param [in] handle The unsampled bindless image handle.
1187+
/// \returns The image_mem_wrapper of unsampled image.
1188+
inline image_mem_wrapper *&get_img_mem(
1189+
const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
1190+
static std::map<sycl::ext::oneapi::experimental::unsampled_image_handle,
1191+
image_mem_wrapper *, detail::sampled_image_handle_compare>
1192+
img_mem_map;
1193+
return img_mem_map[handle];
1194+
}
1195+
1196+
/// Associate an image memory wrapper with a bindless unsampled image handle in
1197+
/// the global registry.
1198+
///
1199+
/// Inserts or updates an entry in the internal image memory map, using the
1200+
/// specified bindless unsampled image handle as the key and the provided image
1201+
/// memory wrapper as the mapped value.
1202+
/// \param [in] handle The unsampled bindless image handle.
1203+
/// \param [in] img_mem The image_mem_wrapper associated with the unsampled handle.
1204+
static inline void set_img_mem(
1205+
const sycl::ext::oneapi::experimental::unsampled_image_handle handle,
1206+
image_mem_wrapper *img_mem) {
1207+
dpct::experimental::get_img_mem(handle) = img_mem;
1208+
}
1209+
1210+
/// Associate an image memory wrapper with a bindless sampled image handle in
1211+
/// the global registry.
1212+
///
1213+
/// Inserts or updates an entry in the internal image memory map, using the
1214+
/// specified bindless unsampled image handle as the key and the provided image
1215+
/// memory wrapper as the mapped value.
1216+
/// \param [in] handle The sampled bindless image handle.
1217+
/// \param [in] img_mem The image_mem_wrapper associated with the sampled handle.
1218+
static inline void set_img_mem(
1219+
const sycl::ext::oneapi::experimental::sampled_image_handle handle,
1220+
image_mem_wrapper *img_mem) {
1221+
dpct::experimental::get_img_mem(handle) = img_mem;
1222+
}
1223+
11881224
#ifdef _WIN32
11891225
/// Map the resource memories to mem handles
11901226
/// \param [in] count The count of resources to map.
@@ -1393,7 +1429,7 @@ create_bindless_image(image_data data, sampling_info info,
13931429
data.get_channel(), data.get_x() / data.get_channel().get_total_size());
13941430
auto img = sycl::ext::oneapi::experimental::create_image(
13951431
mem->get_handle(), samp, mem->get_desc(), q);
1396-
detail::get_img_mem_map(img) = mem;
1432+
set_img_mem(img, mem);
13971433
auto ptr = data.get_data_ptr();
13981434
#ifdef DPCT_USM_LEVEL_NONE
13991435
q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
@@ -1411,7 +1447,7 @@ create_bindless_image(image_data data, sampling_info info,
14111447
new image_mem_wrapper(data.get_channel(), data.get_x(), data.get_y());
14121448
auto img = sycl::ext::oneapi::experimental::create_image(
14131449
mem->get_handle(), samp, mem->get_desc(), q);
1414-
detail::get_img_mem_map(img) = mem;
1450+
set_img_mem(img, mem);
14151451
q.ext_oneapi_copy(
14161452
get_buffer(data.get_data_ptr()).get_host_access().get_pointer(),
14171453
mem->get_handle(), mem->get_desc())
@@ -1455,7 +1491,7 @@ create_bindless_image(image_data data, sycl::queue q = get_default_queue()) {
14551491
data.get_channel(), data.get_x() / data.get_channel().get_total_size());
14561492
auto img = sycl::ext::oneapi::experimental::create_image(
14571493
mem->get_handle(), mem->get_desc(), q);
1458-
detail::get_img_mem_map(img) = mem;
1494+
set_img_mem(img, mem);
14591495
auto ptr = data.get_data_ptr();
14601496
#ifdef DPCT_USM_LEVEL_NONE
14611497
q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
@@ -1471,7 +1507,7 @@ create_bindless_image(image_data data, sycl::queue q = get_default_queue()) {
14711507
new image_mem_wrapper(data.get_channel(), data.get_x(), data.get_y());
14721508
auto img = sycl::ext::oneapi::experimental::create_image(
14731509
mem->get_handle(), mem->get_desc(), q);
1474-
detail::get_img_mem_map(img) = mem;
1510+
set_img_mem(img, mem);
14751511
#ifdef DPCT_USM_LEVEL_NONE
14761512
q.ext_oneapi_copy(
14771513
get_buffer(data.get_data_ptr()).get_host_access().get_pointer(),
@@ -1505,7 +1541,7 @@ create_bindless_image(image_data data, sycl::queue q = get_default_queue()) {
15051541
template <class T>
15061542
static inline void destroy_bindless_image(T handle,
15071543
sycl::queue q = get_default_queue()) {
1508-
auto &mem = detail::get_img_mem_map(handle);
1544+
auto &mem = get_img_mem(handle);
15091545
if (mem) {
15101546
delete mem;
15111547
mem = nullptr;
@@ -1556,7 +1592,7 @@ class bindless_image_wrapper_base {
15561592
auto mem = new image_mem_wrapper(channel, size);
15571593
_img = sycl::ext::oneapi::experimental::create_image(
15581594
mem->get_handle(), samp, mem->get_desc(), q);
1559-
detail::get_img_mem_map(_img) = mem;
1595+
set_img_mem(_img, mem);
15601596
auto ptr = data;
15611597
#ifdef DPCT_USM_LEVEL_NONE
15621598
q.ext_oneapi_copy(get_buffer(data).get_host_access().get_pointer(),
@@ -1590,7 +1626,7 @@ class bindless_image_wrapper_base {
15901626
auto mem = new image_mem_wrapper(desc);
15911627
_img = sycl::ext::oneapi::experimental::create_image(mem->get_handle(),
15921628
samp, *desc, q);
1593-
detail::get_img_mem_map(_img) = mem;
1629+
set_img_mem(_img, mem);
15941630
q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
15951631
mem->get_handle(), mem->get_desc())
15961632
.wait();
@@ -1760,7 +1796,7 @@ class bindless_image_wrapper_base {
17601796
/// Get mipmap memory wrapper attached the bindless image
17611797
/// \return The mipmap memory wrapper
17621798
inline image_mem_wrapper *get_attached_mipmap_data(void) {
1763-
auto mem = detail::get_img_mem_map(_img);
1799+
auto mem = get_img_mem(_img);
17641800

17651801
if (mem->get_image_type() !=
17661802
sycl::ext::oneapi::experimental::image_type::mipmap)

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -180,8 +180,13 @@ void driverMemoryManagement() {
180180
p3d.Depth = s;
181181
// CHECK: *pArr = new dpct::experimental::image_mem_wrapper(&p3DDesc);
182182
cuArray3DCreate(pArr, &p3DDesc);
183+
// CHECK: p3DDesc = (*pArr)->get_desc();
184+
cuArray3DGetDescriptor(&p3DDesc, *pArr);
183185
// CHECK: *pArr = new dpct::experimental::image_mem_wrapper(&pDesc);
184186
cuArrayCreate(pArr, &pDesc);
187+
// CHECK: pDesc = (*pArr)->get_desc();
188+
cuArrayGetDescriptor(&pDesc, *pArr);
189+
185190
// CHECK: delete (*pArr);
186191
cuArrayDestroy(*pArr);
187192
// CHECK: dpct::dpct_memcpy(p2d);

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

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,10 @@ void driverTextureReferenceManagement() {
4747
CUarray a;
4848
// CHECK: sycl::image_channel_type f;
4949
CUarray_format f;
50+
// CHECK: r = new dpct::experimental::bindless_image_wrapper_base();
51+
cuTexRefCreate(&r);
52+
// CHECK: delete r;
53+
cuTexRefDestroy(r);
5054
// CHECK: am = r->get_addressing_mode();
5155
cuTexRefGetAddressMode(&am, r, i);
5256
// CHECK: fm = r->get_filtering_mode();
@@ -72,7 +76,16 @@ void driverTextureReferenceManagement() {
7276
// CHECK-NEXT: r->set_channel_num(i);
7377
cuTexRefSetFormat(r, f, i);
7478
}
75-
79+
void test_surf_ref() {
80+
// CHECK: dpct::experimental::image_mem_wrapper_ptr arr;
81+
CUarray arr;
82+
// CHECK: sycl::ext::oneapi::experimental::unsampled_image_handle ref;
83+
CUsurfref ref;
84+
// CHECK: arr = dpct::experimental::get_img_mem(ref);
85+
cuSurfRefGetArray(&arr, ref);
86+
// CHECK: dpct::experimental::set_img_mem(ref, arr);
87+
cuSurfRefSetArray(ref, arr, 0);
88+
}
7689
int main() {
7790
int i;
7891
// CHECK: tex2.set(sycl::addressing_mode::repeat);

0 commit comments

Comments
 (0)