@@ -2241,6 +2241,25 @@ void handler::setKernelIsCooperative(bool KernelIsCooperative) {
2241
2241
impl->MKernelIsCooperative = KernelIsCooperative;
2242
2242
}
2243
2243
2244
+ #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2245
+ void handler::setKernelClusterLaunch (sycl::range<3 > ClusterSize, int Dims) {
2246
+ throwIfGraphAssociated<
2247
+ syclex::detail::UnsupportedGraphFeatures::
2248
+ sycl_ext_oneapi_experimental_cuda_cluster_launch>();
2249
+ impl->MKernelUsesClusterLaunch = true ;
2250
+
2251
+ if (Dims == 1 ) {
2252
+ sycl::range<1 > ClusterSizeTrimmed = {ClusterSize[0 ]};
2253
+ impl->MNDRDesc .setClusterDimensions (ClusterSizeTrimmed);
2254
+ } else if (Dims == 2 ) {
2255
+ sycl::range<2 > ClusterSizeTrimmed = {ClusterSize[0 ], ClusterSize[1 ]};
2256
+ impl->MNDRDesc .setClusterDimensions (ClusterSizeTrimmed);
2257
+ } else if (Dims == 3 ) {
2258
+ impl->MNDRDesc .setClusterDimensions (ClusterSize);
2259
+ }
2260
+ }
2261
+ #endif
2262
+
2244
2263
void handler::setKernelClusterLaunch (sycl::range<3 > ClusterSize) {
2245
2264
throwIfGraphAssociated<
2246
2265
syclex::detail::UnsupportedGraphFeatures::
@@ -2417,6 +2436,56 @@ bool handler::HasAssociatedAccessor(detail::AccessorImplHost *Req,
2417
2436
void handler::setType (sycl::detail::CGType Type) { impl->MCGType = Type; }
2418
2437
sycl::detail::CGType handler::getType () const { return impl->MCGType ; }
2419
2438
2439
+ #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
2440
+ void handler::setNDRangeDescriptorPadded (sycl::range<3 > N,
2441
+ bool SetNumWorkGroups, int Dims) {
2442
+ if (Dims == 1 ) {
2443
+ sycl::range<1 > Range = {N[0 ]};
2444
+ impl->MNDRDesc = NDRDescT{Range, SetNumWorkGroups};
2445
+ } else if (Dims == 2 ) {
2446
+ sycl::range<2 > Range = {N[0 ], N[1 ]};
2447
+ impl->MNDRDesc = NDRDescT{Range, SetNumWorkGroups};
2448
+ } else if (Dims == 3 ) {
2449
+ impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups};
2450
+ }
2451
+ }
2452
+
2453
+ void handler::setNDRangeDescriptorPadded (sycl::range<3 > NumWorkItems,
2454
+ sycl::id<3 > Offset, int Dims) {
2455
+ if (Dims == 1 ) {
2456
+ sycl::range<1 > NumWorkItemsTrimmed = {NumWorkItems[0 ]};
2457
+ sycl::id<1 > OffsetTrimmed = {Offset[0 ]};
2458
+ impl->MNDRDesc = NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed};
2459
+ } else if (Dims == 2 ) {
2460
+ sycl::range<2 > NumWorkItemsTrimmed = {NumWorkItems[0 ], NumWorkItems[1 ]};
2461
+ sycl::id<2 > OffsetTrimmed = {Offset[0 ], Offset[1 ]};
2462
+ impl->MNDRDesc = NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed};
2463
+ } else if (Dims == 3 ) {
2464
+ impl->MNDRDesc = NDRDescT{NumWorkItems, Offset};
2465
+ }
2466
+ }
2467
+
2468
+ void handler::setNDRangeDescriptorPadded (sycl::range<3 > NumWorkItems,
2469
+ sycl::range<3 > LocalSize,
2470
+ sycl::id<3 > Offset, int Dims) {
2471
+ if (Dims == 1 ) {
2472
+ sycl::range<1 > NumWorkItemsTrimmed = {NumWorkItems[0 ]};
2473
+ sycl::range<1 > LocalSizeTrimmed = {LocalSize[0 ]};
2474
+ sycl::id<1 > OffsetTrimmed = {Offset[0 ]};
2475
+ impl->MNDRDesc =
2476
+ NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed};
2477
+ } else if (Dims == 2 ) {
2478
+ sycl::range<2 > NumWorkItemsTrimmed = {NumWorkItems[0 ], NumWorkItems[1 ]};
2479
+ sycl::range<2 > LocalSizeTrimmed = {LocalSize[0 ], LocalSize[1 ]};
2480
+ sycl::id<2 > OffsetTrimmed = {Offset[0 ], Offset[1 ]};
2481
+ impl->MNDRDesc =
2482
+ NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed};
2483
+ } else if (Dims == 3 ) {
2484
+ impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset};
2485
+ }
2486
+ }
2487
+ #endif
2488
+
2420
2489
void handler::setNDRangeDescriptor (sycl::range<3 > N, bool SetNumWorkGroups) {
2421
2490
impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups};
2422
2491
}
0 commit comments