16 inline namespace _V1 {
17 namespace ext::oneapi::experimental {
20 #define __SYCL_ARCHITECTURE(NAME, VAL) NAME = VAL,
21 #define __SYCL_ARCHITECTURE_ALIAS(NAME, VAL) NAME = VAL,
22 #include <sycl/ext/oneapi/experimental/architectures.def>
23 #undef __SYCL_ARCHITECTURE
24 #undef __SYCL_ARCHITECTURE_ALIAS
52 ext::oneapi::experimental::architecture::intel_gpu_bdw;
55 ext::oneapi::experimental::architecture::intel_gpu_lnl_m;
59 ext::oneapi::experimental::architecture::nvidia_gpu_sm_50;
62 ext::oneapi::experimental::architecture::nvidia_gpu_sm_90a;
66 ext::oneapi::experimental::architecture::amd_gpu_gfx700;
69 ext::oneapi::experimental::architecture::amd_gpu_gfx1201;
71 #ifndef __SYCL_TARGET_INTEL_X86_64__
72 #define __SYCL_TARGET_INTEL_X86_64__ 0
74 #ifndef __SYCL_TARGET_INTEL_GPU_BDW__
75 #define __SYCL_TARGET_INTEL_GPU_BDW__ 0
77 #ifndef __SYCL_TARGET_INTEL_GPU_SKL__
78 #define __SYCL_TARGET_INTEL_GPU_SKL__ 0
80 #ifndef __SYCL_TARGET_INTEL_GPU_KBL__
81 #define __SYCL_TARGET_INTEL_GPU_KBL__ 0
83 #ifndef __SYCL_TARGET_INTEL_GPU_CFL__
84 #define __SYCL_TARGET_INTEL_GPU_CFL__ 0
86 #ifndef __SYCL_TARGET_INTEL_GPU_APL__
87 #define __SYCL_TARGET_INTEL_GPU_APL__ 0
89 #ifndef __SYCL_TARGET_INTEL_GPU_GLK__
90 #define __SYCL_TARGET_INTEL_GPU_GLK__ 0
92 #ifndef __SYCL_TARGET_INTEL_GPU_WHL__
93 #define __SYCL_TARGET_INTEL_GPU_WHL__ 0
95 #ifndef __SYCL_TARGET_INTEL_GPU_AML__
96 #define __SYCL_TARGET_INTEL_GPU_AML__ 0
98 #ifndef __SYCL_TARGET_INTEL_GPU_CML__
99 #define __SYCL_TARGET_INTEL_GPU_CML__ 0
101 #ifndef __SYCL_TARGET_INTEL_GPU_ICLLP__
102 #define __SYCL_TARGET_INTEL_GPU_ICLLP__ 0
104 #ifndef __SYCL_TARGET_INTEL_GPU_EHL__
105 #define __SYCL_TARGET_INTEL_GPU_EHL__ 0
107 #ifndef __SYCL_TARGET_INTEL_GPU_TGLLP__
108 #define __SYCL_TARGET_INTEL_GPU_TGLLP__ 0
110 #ifndef __SYCL_TARGET_INTEL_GPU_RKL__
111 #define __SYCL_TARGET_INTEL_GPU_RKL__ 0
113 #ifndef __SYCL_TARGET_INTEL_GPU_ADL_S__
114 #define __SYCL_TARGET_INTEL_GPU_ADL_S__ 0
116 #ifndef __SYCL_TARGET_INTEL_GPU_ADL_P__
117 #define __SYCL_TARGET_INTEL_GPU_ADL_P__ 0
119 #ifndef __SYCL_TARGET_INTEL_GPU_ADL_N__
120 #define __SYCL_TARGET_INTEL_GPU_ADL_N__ 0
122 #ifndef __SYCL_TARGET_INTEL_GPU_DG1__
123 #define __SYCL_TARGET_INTEL_GPU_DG1__ 0
125 #ifndef __SYCL_TARGET_INTEL_GPU_ACM_G10__
126 #define __SYCL_TARGET_INTEL_GPU_ACM_G10__ 0
128 #ifndef __SYCL_TARGET_INTEL_GPU_ACM_G11__
129 #define __SYCL_TARGET_INTEL_GPU_ACM_G11__ 0
131 #ifndef __SYCL_TARGET_INTEL_GPU_ACM_G12__
132 #define __SYCL_TARGET_INTEL_GPU_ACM_G12__ 0
134 #ifndef __SYCL_TARGET_INTEL_GPU_PVC__
135 #define __SYCL_TARGET_INTEL_GPU_PVC__ 0
137 #ifndef __SYCL_TARGET_INTEL_GPU_PVC_VG__
138 #define __SYCL_TARGET_INTEL_GPU_PVC_VG__ 0
140 #ifndef __SYCL_TARGET_INTEL_GPU_MTL_U__
141 #define __SYCL_TARGET_INTEL_GPU_MTL_U__ 0
143 #ifndef __SYCL_TARGET_INTEL_GPU_MTL_H__
144 #define __SYCL_TARGET_INTEL_GPU_MTL_H__ 0
146 #ifndef __SYCL_TARGET_INTEL_GPU_ARL_H__
147 #define __SYCL_TARGET_INTEL_GPU_ARL_H__ 0
149 #ifndef __SYCL_TARGET_INTEL_GPU_BMG_G21__
150 #define __SYCL_TARGET_INTEL_GPU_BMG_G21__ 0
152 #ifndef __SYCL_TARGET_INTEL_GPU_LNL_M__
153 #define __SYCL_TARGET_INTEL_GPU_LNL_M__ 0
155 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM50__
156 #define __SYCL_TARGET_NVIDIA_GPU_SM50__ 0
158 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM52__
159 #define __SYCL_TARGET_NVIDIA_GPU_SM52__ 0
161 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM53__
162 #define __SYCL_TARGET_NVIDIA_GPU_SM53__ 0
164 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM60__
165 #define __SYCL_TARGET_NVIDIA_GPU_SM60__ 0
167 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM61__
168 #define __SYCL_TARGET_NVIDIA_GPU_SM61__ 0
170 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM62__
171 #define __SYCL_TARGET_NVIDIA_GPU_SM62__ 0
173 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM70__
174 #define __SYCL_TARGET_NVIDIA_GPU_SM70__ 0
176 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM72__
177 #define __SYCL_TARGET_NVIDIA_GPU_SM72__ 0
179 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM75__
180 #define __SYCL_TARGET_NVIDIA_GPU_SM75__ 0
182 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM80__
183 #define __SYCL_TARGET_NVIDIA_GPU_SM80__ 0
185 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM86__
186 #define __SYCL_TARGET_NVIDIA_GPU_SM86__ 0
188 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM87__
189 #define __SYCL_TARGET_NVIDIA_GPU_SM87__ 0
191 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM89__
192 #define __SYCL_TARGET_NVIDIA_GPU_SM89__ 0
194 #ifndef __SYCL_TARGET_NVIDIA_GPU_SM90__
195 #define __SYCL_TARGET_NVIDIA_GPU_SM90__ 0
197 #ifndef __SYCL_TARGET_AMD_GPU_GFX700__
198 #define __SYCL_TARGET_AMD_GPU_GFX700__ 0
200 #ifndef __SYCL_TARGET_AMD_GPU_GFX701__
201 #define __SYCL_TARGET_AMD_GPU_GFX701__ 0
203 #ifndef __SYCL_TARGET_AMD_GPU_GFX702__
204 #define __SYCL_TARGET_AMD_GPU_GFX702__ 0
206 #ifndef __SYCL_TARGET_AMD_GPU_GFX801__
207 #define __SYCL_TARGET_AMD_GPU_GFX801__ 0
209 #ifndef __SYCL_TARGET_AMD_GPU_GFX802__
210 #define __SYCL_TARGET_AMD_GPU_GFX802__ 0
212 #ifndef __SYCL_TARGET_AMD_GPU_GFX803__
213 #define __SYCL_TARGET_AMD_GPU_GFX803__ 0
215 #ifndef __SYCL_TARGET_AMD_GPU_GFX805__
216 #define __SYCL_TARGET_AMD_GPU_GFX805__ 0
218 #ifndef __SYCL_TARGET_AMD_GPU_GFX810__
219 #define __SYCL_TARGET_AMD_GPU_GFX810__ 0
221 #ifndef __SYCL_TARGET_AMD_GPU_GFX900__
222 #define __SYCL_TARGET_AMD_GPU_GFX900__ 0
224 #ifndef __SYCL_TARGET_AMD_GPU_GFX902__
225 #define __SYCL_TARGET_AMD_GPU_GFX902__ 0
227 #ifndef __SYCL_TARGET_AMD_GPU_GFX904__
228 #define __SYCL_TARGET_AMD_GPU_GFX904__ 0
230 #ifndef __SYCL_TARGET_AMD_GPU_GFX906__
231 #define __SYCL_TARGET_AMD_GPU_GFX906__ 0
233 #ifndef __SYCL_TARGET_AMD_GPU_GFX908__
234 #define __SYCL_TARGET_AMD_GPU_GFX908__ 0
236 #ifndef __SYCL_TARGET_AMD_GPU_GFX909__
237 #define __SYCL_TARGET_AMD_GPU_GFX909__ 0
239 #ifndef __SYCL_TARGET_AMD_GPU_GFX90A__
240 #define __SYCL_TARGET_AMD_GPU_GFX90A__ 0
242 #ifndef __SYCL_TARGET_AMD_GPU_GFX90C__
243 #define __SYCL_TARGET_AMD_GPU_GFX90C__ 0
245 #ifndef __SYCL_TARGET_AMD_GPU_GFX940__
246 #define __SYCL_TARGET_AMD_GPU_GFX940__ 0
248 #ifndef __SYCL_TARGET_AMD_GPU_GFX941__
249 #define __SYCL_TARGET_AMD_GPU_GFX941__ 0
251 #ifndef __SYCL_TARGET_AMD_GPU_GFX942__
252 #define __SYCL_TARGET_AMD_GPU_GFX942__ 0
254 #ifndef __SYCL_TARGET_AMD_GPU_GFX1010__
255 #define __SYCL_TARGET_AMD_GPU_GFX1010__ 0
257 #ifndef __SYCL_TARGET_AMD_GPU_GFX1011__
258 #define __SYCL_TARGET_AMD_GPU_GFX1011__ 0
260 #ifndef __SYCL_TARGET_AMD_GPU_GFX1012__
261 #define __SYCL_TARGET_AMD_GPU_GFX1012__ 0
263 #ifndef __SYCL_TARGET_AMD_GPU_GFX1013__
264 #define __SYCL_TARGET_AMD_GPU_GFX1013__ 0
266 #ifndef __SYCL_TARGET_AMD_GPU_GFX1030__
267 #define __SYCL_TARGET_AMD_GPU_GFX1030__ 0
269 #ifndef __SYCL_TARGET_AMD_GPU_GFX1031__
270 #define __SYCL_TARGET_AMD_GPU_GFX1031__ 0
272 #ifndef __SYCL_TARGET_AMD_GPU_GFX1032__
273 #define __SYCL_TARGET_AMD_GPU_GFX1032__ 0
275 #ifndef __SYCL_TARGET_AMD_GPU_GFX1033__
276 #define __SYCL_TARGET_AMD_GPU_GFX1033__ 0
278 #ifndef __SYCL_TARGET_AMD_GPU_GFX1034__
279 #define __SYCL_TARGET_AMD_GPU_GFX1034__ 0
281 #ifndef __SYCL_TARGET_AMD_GPU_GFX1035__
282 #define __SYCL_TARGET_AMD_GPU_GFX1035__ 0
284 #ifndef __SYCL_TARGET_AMD_GPU_GFX1036__
285 #define __SYCL_TARGET_AMD_GPU_GFX1036__ 0
287 #ifndef __SYCL_TARGET_AMD_GPU_GFX1100__
288 #define __SYCL_TARGET_AMD_GPU_GFX1100__ 0
290 #ifndef __SYCL_TARGET_AMD_GPU_GFX1101__
291 #define __SYCL_TARGET_AMD_GPU_GFX1101__ 0
293 #ifndef __SYCL_TARGET_AMD_GPU_GFX1102__
294 #define __SYCL_TARGET_AMD_GPU_GFX1102__ 0
296 #ifndef __SYCL_TARGET_AMD_GPU_GFX1103__
297 #define __SYCL_TARGET_AMD_GPU_GFX1103__ 0
299 #ifndef __SYCL_TARGET_AMD_GPU_GFX1150__
300 #define __SYCL_TARGET_AMD_GPU_GFX1150__ 0
302 #ifndef __SYCL_TARGET_AMD_GPU_GFX1151__
303 #define __SYCL_TARGET_AMD_GPU_GFX1151__ 0
305 #ifndef __SYCL_TARGET_AMD_GPU_GFX1200__
306 #define __SYCL_TARGET_AMD_GPU_GFX1200__ 0
308 #ifndef __SYCL_TARGET_AMD_GPU_GFX1201__
309 #define __SYCL_TARGET_AMD_GPU_GFX1201__ 0
399 constexpr
static std::optional<ext::oneapi::experimental::architecture>
403 #if __SYCL_TARGET_INTEL_X86_64__
404 return ext::oneapi::experimental::architecture::x86_64;
406 #if __SYCL_TARGET_INTEL_GPU_BDW__
407 return ext::oneapi::experimental::architecture::intel_gpu_bdw;
409 #if __SYCL_TARGET_INTEL_GPU_SKL__
410 return ext::oneapi::experimental::architecture::intel_gpu_skl;
412 #if __SYCL_TARGET_INTEL_GPU_KBL__
413 return ext::oneapi::experimental::architecture::intel_gpu_kbl;
415 #if __SYCL_TARGET_INTEL_GPU_CFL__
416 return ext::oneapi::experimental::architecture::intel_gpu_cfl;
418 #if __SYCL_TARGET_INTEL_GPU_APL__
419 return ext::oneapi::experimental::architecture::intel_gpu_apl;
421 #if __SYCL_TARGET_INTEL_GPU_GLK__
422 return ext::oneapi::experimental::architecture::intel_gpu_glk;
424 #if __SYCL_TARGET_INTEL_GPU_WHL__
425 return ext::oneapi::experimental::architecture::intel_gpu_whl;
427 #if __SYCL_TARGET_INTEL_GPU_AML__
428 return ext::oneapi::experimental::architecture::intel_gpu_aml;
430 #if __SYCL_TARGET_INTEL_GPU_CML__
431 return ext::oneapi::experimental::architecture::intel_gpu_cml;
433 #if __SYCL_TARGET_INTEL_GPU_ICLLP__
434 return ext::oneapi::experimental::architecture::intel_gpu_icllp;
436 #if __SYCL_TARGET_INTEL_GPU_EHL__
437 return ext::oneapi::experimental::architecture::intel_gpu_ehl;
439 #if __SYCL_TARGET_INTEL_GPU_TGLLP__
440 return ext::oneapi::experimental::architecture::intel_gpu_tgllp;
442 #if __SYCL_TARGET_INTEL_GPU_RKL__
443 return ext::oneapi::experimental::architecture::intel_gpu_rkl;
445 #if __SYCL_TARGET_INTEL_GPU_ADL_S__
446 return ext::oneapi::experimental::architecture::intel_gpu_adl_s;
448 #if __SYCL_TARGET_INTEL_GPU_ADL_P__
449 return ext::oneapi::experimental::architecture::intel_gpu_adl_p;
451 #if __SYCL_TARGET_INTEL_GPU_ADL_P__
452 return ext::oneapi::experimental::architecture::intel_gpu_adl_p;
454 #if __SYCL_TARGET_INTEL_GPU_ADL_N__
455 return ext::oneapi::experimental::architecture::intel_gpu_adl_n;
457 #if __SYCL_TARGET_INTEL_GPU_DG1__
458 return ext::oneapi::experimental::architecture::intel_gpu_dg1;
460 #if __SYCL_TARGET_INTEL_GPU_ACM_G10__
461 return ext::oneapi::experimental::architecture::intel_gpu_acm_g10;
463 #if __SYCL_TARGET_INTEL_GPU_ACM_G11__
464 return ext::oneapi::experimental::architecture::intel_gpu_acm_g11;
466 #if __SYCL_TARGET_INTEL_GPU_ACM_G12__
467 return ext::oneapi::experimental::architecture::intel_gpu_acm_g12;
469 #if __SYCL_TARGET_INTEL_GPU_PVC__
470 return ext::oneapi::experimental::architecture::intel_gpu_pvc;
472 #if __SYCL_TARGET_INTEL_GPU_PVC_VG__
473 return ext::oneapi::experimental::architecture::intel_gpu_pvc_vg;
475 #if __SYCL_TARGET_INTEL_GPU_MTL_U__
476 return ext::oneapi::experimental::architecture::intel_gpu_mtl_u;
478 #if __SYCL_TARGET_INTEL_GPU_MTL_H__
479 return ext::oneapi::experimental::architecture::intel_gpu_mtl_h;
481 #if __SYCL_TARGET_INTEL_GPU_ARL_H__
482 return ext::oneapi::experimental::architecture::intel_gpu_arl_h;
484 #if __SYCL_TARGET_INTEL_GPU_BMG_G21__
485 return ext::oneapi::experimental::architecture::intel_gpu_bmg_g21;
487 #if __SYCL_TARGET_INTEL_GPU_LNL_M__
488 return ext::oneapi::experimental::architecture::intel_gpu_lnl_m;
490 #if __SYCL_TARGET_NVIDIA_GPU_SM50__
491 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_50;
493 #if __SYCL_TARGET_NVIDIA_GPU_SM52__
494 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_52;
496 #if __SYCL_TARGET_NVIDIA_GPU_SM53__
497 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_53;
499 #if __SYCL_TARGET_NVIDIA_GPU_SM60__
500 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_60;
502 #if __SYCL_TARGET_NVIDIA_GPU_SM61__
503 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_61;
505 #if __SYCL_TARGET_NVIDIA_GPU_SM62__
506 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_62;
508 #if __SYCL_TARGET_NVIDIA_GPU_SM70__
509 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_70;
511 #if __SYCL_TARGET_NVIDIA_GPU_SM72__
512 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_72;
514 #if __SYCL_TARGET_NVIDIA_GPU_SM75__
515 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_75;
517 #if __SYCL_TARGET_NVIDIA_GPU_SM80__
518 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_80;
520 #if __SYCL_TARGET_NVIDIA_GPU_SM86__
521 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_86;
523 #if __SYCL_TARGET_NVIDIA_GPU_SM87__
524 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_87;
526 #if __SYCL_TARGET_NVIDIA_GPU_SM89__
527 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_89;
529 #if __SYCL_TARGET_NVIDIA_GPU_SM90__
530 return ext::oneapi::experimental::architecture::nvidia_gpu_sm_90;
532 #if __SYCL_TARGET_AMD_GPU_GFX700__
533 return ext::oneapi::experimental::architecture::amd_gpu_gfx700;
535 #if __SYCL_TARGET_AMD_GPU_GFX701__
536 return ext::oneapi::experimental::architecture::amd_gpu_gfx701;
538 #if __SYCL_TARGET_AMD_GPU_GFX702__
539 return ext::oneapi::experimental::architecture::amd_gpu_gfx702;
541 #if __SYCL_TARGET_AMD_GPU_GFX801__
542 return ext::oneapi::experimental::architecture::amd_gpu_gfx801;
544 #if __SYCL_TARGET_AMD_GPU_GFX802__
545 return ext::oneapi::experimental::architecture::amd_gpu_gfx802;
547 #if __SYCL_TARGET_AMD_GPU_GFX803__
548 return ext::oneapi::experimental::architecture::amd_gpu_gfx803;
550 #if __SYCL_TARGET_AMD_GPU_GFX805__
551 return ext::oneapi::experimental::architecture::amd_gpu_gfx805;
553 #if __SYCL_TARGET_AMD_GPU_GFX810__
554 return ext::oneapi::experimental::architecture::amd_gpu_gfx810;
556 #if __SYCL_TARGET_AMD_GPU_GFX900__
557 return ext::oneapi::experimental::architecture::amd_gpu_gfx900;
559 #if __SYCL_TARGET_AMD_GPU_GFX902__
560 return ext::oneapi::experimental::architecture::amd_gpu_gfx902;
562 #if __SYCL_TARGET_AMD_GPU_GFX904__
563 return ext::oneapi::experimental::architecture::amd_gpu_gfx904;
565 #if __SYCL_TARGET_AMD_GPU_GFX906__
566 return ext::oneapi::experimental::architecture::amd_gpu_gfx906;
568 #if __SYCL_TARGET_AMD_GPU_GFX908__
569 return ext::oneapi::experimental::architecture::amd_gpu_gfx908;
571 #if __SYCL_TARGET_AMD_GPU_GFX909__
572 return ext::oneapi::experimental::architecture::amd_gpu_gfx909;
574 #if __SYCL_TARGET_AMD_GPU_GFX90a__
575 return ext::oneapi::experimental::architecture::amd_gpu_gfx90a;
577 #if __SYCL_TARGET_AMD_GPU_GFX90c__
578 return ext::oneapi::experimental::architecture::amd_gpu_gfx90c;
580 #if __SYCL_TARGET_AMD_GPU_GFX940__
581 return ext::oneapi::experimental::architecture::amd_gpu_gfx940;
583 #if __SYCL_TARGET_AMD_GPU_GFX941__
584 return ext::oneapi::experimental::architecture::amd_gpu_gfx941;
586 #if __SYCL_TARGET_AMD_GPU_GFX942__
587 return ext::oneapi::experimental::architecture::amd_gpu_gfx942;
589 #if __SYCL_TARGET_AMD_GPU_GFX1010__
590 return ext::oneapi::experimental::architecture::amd_gpu_gfx1010;
592 #if __SYCL_TARGET_AMD_GPU_GFX1011__
593 return ext::oneapi::experimental::architecture::amd_gpu_gfx1011;
595 #if __SYCL_TARGET_AMD_GPU_GFX1012__
596 return ext::oneapi::experimental::architecture::amd_gpu_gfx1012;
598 #if __SYCL_TARGET_AMD_GPU_GFX1030__
599 return ext::oneapi::experimental::architecture::amd_gpu_gfx1030;
601 #if __SYCL_TARGET_AMD_GPU_GFX1031__
602 return ext::oneapi::experimental::architecture::amd_gpu_gfx1031;
604 #if __SYCL_TARGET_AMD_GPU_GFX1032__
605 return ext::oneapi::experimental::architecture::amd_gpu_gfx1032;
607 #if __SYCL_TARGET_AMD_GPU_GFX1033__
608 return ext::oneapi::experimental::architecture::amd_gpu_gfx1033;
610 #if __SYCL_TARGET_AMD_GPU_GFX1034__
611 return ext::oneapi::experimental::architecture::amd_gpu_gfx1034;
613 #if __SYCL_TARGET_AMD_GPU_GFX1035__
614 return ext::oneapi::experimental::architecture::amd_gpu_gfx1035;
616 #if __SYCL_TARGET_AMD_GPU_GFX1036__
617 return ext::oneapi::experimental::architecture::amd_gpu_gfx1036;
619 #if __SYCL_TARGET_AMD_GPU_GFX1100__
620 return ext::oneapi::experimental::architecture::amd_gpu_gfx1100;
622 #if __SYCL_TARGET_AMD_GPU_GFX1101__
623 return ext::oneapi::experimental::architecture::amd_gpu_gfx1101;
625 #if __SYCL_TARGET_AMD_GPU_GFX1102__
626 return ext::oneapi::experimental::architecture::amd_gpu_gfx1102;
628 #if __SYCL_TARGET_AMD_GPU_GFX1103__
629 return ext::oneapi::experimental::architecture::amd_gpu_gfx1103;
631 #if __SYCL_TARGET_AMD_GPU_GFX1150__
632 return ext::oneapi::experimental::architecture::amd_gpu_gfx1150;
634 #if __SYCL_TARGET_AMD_GPU_GFX1151__
635 return ext::oneapi::experimental::architecture::amd_gpu_gfx1151;
637 #if __SYCL_TARGET_AMD_GPU_GFX1200__
638 return ext::oneapi::experimental::architecture::amd_gpu_gfx1200;
640 #if __SYCL_TARGET_AMD_GPU_GFX1201__
641 return ext::oneapi::experimental::architecture::amd_gpu_gfx1201;
647 constexpr
static bool
649 constexpr std::optional<ext::oneapi::experimental::architecture>
651 if (current_arch.has_value())
652 return arch == *current_arch;
670 static constexpr std::optional<ext::oneapi::experimental::architecture>
683 static constexpr std::optional<ext::oneapi::experimental::architecture>
696 template <ext::oneapi::experimental::arch_category Category>
698 constexpr std::optional<ext::oneapi::experimental::architecture>
700 constexpr std::optional<ext::oneapi::experimental::architecture>
702 constexpr std::optional<ext::oneapi::experimental::architecture>
705 if (category_min_arch.has_value() && category_max_arch.has_value() &&
706 current_arch.has_value())
707 if ((*category_min_arch <= *current_arch) &&
708 (*current_arch <= *category_max_arch))
716 return (device_architecture_is_in_category_aot<Categories>() || ...);
719 constexpr
static std::optional<ext::oneapi::experimental::arch_category>
721 auto arch_is_in_segment =
724 if ((
min <= arch) && (arch <=
max))
742 template <ext::oneapi::experimental::architecture Arch,
typename Compare>
744 constexpr std::optional<ext::oneapi::experimental::arch_category>
746 constexpr std::optional<ext::oneapi::experimental::architecture>
749 if (input_arch_category.has_value() && current_arch.has_value()) {
750 std::optional<ext::oneapi::experimental::arch_category>
752 if (current_arch_category.has_value() &&
753 (*input_arch_category == *current_arch_category))
754 return comp(*current_arch, Arch);
786 if constexpr (MakeCall && device_architecture_is<Archs...>()) {
801 if constexpr (MakeCall &&
802 device_architecture_is_in_categories<Categories...>()) {
815 template <ext::oneapi::experimental::architecture Arch,
typename T>
817 if constexpr (MakeCall &&
818 sycl::detail::device_architecture_comparison_aot<Arch>(
821 return sycl::detail::if_architecture_helper<false>{};
824 return sycl::detail::if_architecture_helper<MakeCall>{};
833 template <ext::oneapi::experimental::architecture Arch,
typename T>
835 if constexpr (MakeCall &&
836 sycl::detail::device_architecture_comparison_aot<Arch>(
839 return sycl::detail::if_architecture_helper<false>{};
842 return sycl::detail::if_architecture_helper<MakeCall>{};
850 template <ext::oneapi::experimental::architecture Arch,
typename T>
852 if constexpr (MakeCall &&
853 sycl::detail::device_architecture_comparison_aot<Arch>(
856 return sycl::detail::if_architecture_helper<false>{};
859 return sycl::detail::if_architecture_helper<MakeCall>{};
868 template <ext::oneapi::experimental::architecture Arch,
typename T>
870 if constexpr (MakeCall &&
871 sycl::detail::device_architecture_comparison_aot<Arch>(
874 return sycl::detail::if_architecture_helper<false>{};
877 return sycl::detail::if_architecture_helper<MakeCall>{};
889 if constexpr (MakeCall &&
890 sycl::detail::device_architecture_comparison_aot<Arch1>(
892 sycl::detail::device_architecture_comparison_aot<Arch2>(
895 return sycl::detail::if_architecture_helper<false>{};
898 return sycl::detail::if_architecture_helper<MakeCall>{};
903 if constexpr (MakeCall) {
910 namespace ext::oneapi::experimental {
919 template <
typename T,
typename... Condition>
920 #ifdef __SYCL_DEVICE_ONLY__
921 [[__sycl_detail__::add_ir_attributes_function(
922 "sycl-call-if-on-device-conditionally",
true)]]
968 template <
int I1,
int I2,
int I3,
int... Is>
969 static auto append(std::integer_sequence<int, Is...>) {
970 return std::integer_sequence<int, I1, I2, I3, Is...>{};
973 static constexpr
int arch =
static_cast<int>(Arch);
975 decltype(append<ConditionOr, ConditionArchitecture, arch>(
rest{}));
979 static constexpr
int arch =
static_cast<int>(Arch);
980 using seq = std::integer_sequence<int, ConditionArchitecture, arch>;
987 template <
int I,
int... Is>
988 static auto append(std::integer_sequence<int, Is...>) {
989 return std::integer_sequence<int, I, Is...>{};
991 using rest =
typename Exp::seq;
992 using seq = decltype(append<ConditionNot>(
rest{}));
999 template <
int I,
int... I1s,
int... I2s>
1000 static auto append(std::integer_sequence<int, I1s...>,
1001 std::integer_sequence<int, I2s...>) {
1002 return std::integer_sequence<int, I, I1s..., I2s...>{};
1013 template <
int I,
int... I1s,
int... I2s>
1014 static auto append(std::integer_sequence<int, I1s...>,
1015 std::integer_sequence<int, I2s...>) {
1016 return std::integer_sequence<int, I, I1s..., I2s...>{};
1026 template <
typename T,
int... Is>
1028 std::integer_sequence<int, Is...>) {
1036 typename = std::enable_if<std::is_invocable_v<T>>>
1038 using make_call_if =
1045 using cond =
typename make_call_if::seq;
1051 using cond =
typename MakeCallIf::seq;
1058 #ifdef SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE_NEW_DESIGN_IMPL
1061 using make_call_if = detail::ConditionAnyArchitectureBuilder<Archs...>;
1062 using make_else_call_if = detail::ConditionNotBuilder<make_call_if>;
1064 using cond =
typename make_call_if::seq;
1066 return detail::if_architecture_is_helper<make_else_call_if>{};
1074 static_assert(sycl::detail::allowable_aot_mode<Archs...>(),
1075 "The if_architecture_is function may only be used when AOT "
1076 "compiling with '-fsycl-targets=spir64_x86_64' or "
1077 "'-fsycl-targets=*_gpu_*'");
1078 if constexpr (sycl::detail::device_architecture_is<Archs...>()) {
1080 return sycl::detail::if_architecture_helper<false>{};
1083 return sycl::detail::if_architecture_helper<true>{};
1096 return sycl::detail::if_architecture_helper<false>{};
1099 return sycl::detail::if_architecture_helper<true>{};
1106 template <architecture Arch,
typename T>
1108 if constexpr (sycl::detail::device_architecture_comparison_aot<Arch>(
1111 return sycl::detail::if_architecture_helper<false>{};
1114 return sycl::detail::if_architecture_helper<true>{};
1121 template <architecture Arch,
typename T>
1123 if constexpr (sycl::detail::device_architecture_comparison_aot<Arch>(
1126 return sycl::detail::if_architecture_helper<false>{};
1129 return sycl::detail::if_architecture_helper<true>{};
1136 template <architecture Arch,
typename T>
1138 if constexpr (sycl::detail::device_architecture_comparison_aot<Arch>(
1141 return sycl::detail::if_architecture_helper<false>{};
1144 return sycl::detail::if_architecture_helper<true>{};
1151 template <architecture Arch,
typename T>
1153 if constexpr (sycl::detail::device_architecture_comparison_aot<Arch>(
1156 return sycl::detail::if_architecture_helper<false>{};
1159 return sycl::detail::if_architecture_helper<true>{};
1167 template <architecture Arch1, architecture Arch2,
typename T>
1169 if constexpr (sycl::detail::device_architecture_comparison_aot<Arch1>(
1171 sycl::detail::device_architecture_comparison_aot<Arch2>(
1174 return sycl::detail::if_architecture_helper<false>{};
1177 return sycl::detail::if_architecture_helper<true>{};
constexpr auto else_if_architecture_is_lt(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
constexpr auto else_if_architecture_is_gt(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
constexpr auto else_if_architecture_is_le(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
constexpr void otherwise(T fn)
constexpr auto else_if_architecture_is(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
constexpr auto else_if_architecture_is_between(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
constexpr auto else_if_architecture_is_ge(T fn)
The condition is true only if the object F comes from a previous call whose associated condition is f...
auto else_if_architecture_is(T fn)
#define __SYCL_TARGET_INTEL_X86_64__
#define __SYCL_TARGET_INTEL_GPU_PVC_VG__
#define __SYCL_TARGET_AMD_GPU_GFX1102__
#define __SYCL_TARGET_AMD_GPU_GFX803__
#define __SYCL_TARGET_INTEL_GPU_GLK__
#define __SYCL_TARGET_AMD_GPU_GFX810__
#define __SYCL_TARGET_NVIDIA_GPU_SM61__
#define __SYCL_TARGET_INTEL_GPU_WHL__
#define __SYCL_TARGET_AMD_GPU_GFX802__
#define __SYCL_TARGET_AMD_GPU_GFX1100__
#define __SYCL_TARGET_NVIDIA_GPU_SM53__
#define __SYCL_TARGET_AMD_GPU_GFX90C__
#define __SYCL_TARGET_INTEL_GPU_CFL__
#define __SYCL_TARGET_AMD_GPU_GFX908__
#define __SYCL_TARGET_AMD_GPU_GFX90A__
#define __SYCL_TARGET_AMD_GPU_GFX1010__
#define __SYCL_TARGET_INTEL_GPU_DG1__
#define __SYCL_TARGET_AMD_GPU_GFX1031__
#define __SYCL_TARGET_AMD_GPU_GFX902__
#define __SYCL_TARGET_AMD_GPU_GFX1103__
#define __SYCL_TARGET_NVIDIA_GPU_SM89__
#define __SYCL_TARGET_INTEL_GPU_ACM_G11__
#define __SYCL_TARGET_INTEL_GPU_ACM_G10__
#define __SYCL_TARGET_INTEL_GPU_LNL_M__
#define __SYCL_TARGET_INTEL_GPU_PVC__
#define __SYCL_TARGET_AMD_GPU_GFX1033__
#define __SYCL_TARGET_AMD_GPU_GFX805__
#define __SYCL_TARGET_AMD_GPU_GFX1201__
#define __SYCL_TARGET_AMD_GPU_GFX1035__
#define __SYCL_TARGET_INTEL_GPU_MTL_H__
#define __SYCL_TARGET_AMD_GPU_GFX942__
#define __SYCL_TARGET_AMD_GPU_GFX1200__
#define __SYCL_TARGET_AMD_GPU_GFX1012__
#define __SYCL_TARGET_INTEL_GPU_KBL__
#define __SYCL_TARGET_AMD_GPU_GFX1151__
#define __SYCL_TARGET_INTEL_GPU_BMG_G21__
#define __SYCL_TARGET_AMD_GPU_GFX1101__
#define __SYCL_TARGET_INTEL_GPU_SKL__
#define __SYCL_TARGET_NVIDIA_GPU_SM50__
#define __SYCL_TARGET_AMD_GPU_GFX1011__
#define __SYCL_TARGET_AMD_GPU_GFX1030__
#define __SYCL_TARGET_AMD_GPU_GFX900__
#define __SYCL_TARGET_INTEL_GPU_ADL_S__
#define __SYCL_TARGET_INTEL_GPU_ADL_N__
#define __SYCL_TARGET_INTEL_GPU_EHL__
#define __SYCL_TARGET_AMD_GPU_GFX904__
#define __SYCL_TARGET_NVIDIA_GPU_SM52__
#define __SYCL_TARGET_AMD_GPU_GFX941__
#define __SYCL_TARGET_AMD_GPU_GFX1036__
#define __SYCL_TARGET_INTEL_GPU_ADL_P__
#define __SYCL_TARGET_AMD_GPU_GFX801__
#define __SYCL_TARGET_NVIDIA_GPU_SM80__
#define __SYCL_TARGET_INTEL_GPU_TGLLP__
#define __SYCL_TARGET_INTEL_GPU_RKL__
#define __SYCL_TARGET_NVIDIA_GPU_SM90__
#define __SYCL_TARGET_INTEL_GPU_AML__
#define __SYCL_TARGET_INTEL_GPU_ICLLP__
#define __SYCL_TARGET_NVIDIA_GPU_SM72__
#define __SYCL_TARGET_AMD_GPU_GFX906__
#define __SYCL_TARGET_INTEL_GPU_BDW__
#define __SYCL_TARGET_NVIDIA_GPU_SM86__
#define __SYCL_TARGET_AMD_GPU_GFX702__
#define __SYCL_TARGET_NVIDIA_GPU_SM60__
#define __SYCL_TARGET_INTEL_GPU_ACM_G12__
#define __SYCL_TARGET_AMD_GPU_GFX1013__
#define __SYCL_TARGET_AMD_GPU_GFX940__
#define __SYCL_TARGET_AMD_GPU_GFX1150__
#define __SYCL_TARGET_NVIDIA_GPU_SM62__
#define __SYCL_TARGET_INTEL_GPU_APL__
#define __SYCL_TARGET_AMD_GPU_GFX701__
#define __SYCL_TARGET_INTEL_GPU_ARL_H__
#define __SYCL_TARGET_AMD_GPU_GFX1034__
#define __SYCL_TARGET_INTEL_GPU_CML__
#define __SYCL_TARGET_INTEL_GPU_MTL_U__
#define __SYCL_TARGET_NVIDIA_GPU_SM70__
#define __SYCL_TARGET_AMD_GPU_GFX700__
#define __SYCL_TARGET_AMD_GPU_GFX1032__
#define __SYCL_TARGET_AMD_GPU_GFX909__
#define __SYCL_TARGET_NVIDIA_GPU_SM87__
#define __SYCL_TARGET_NVIDIA_GPU_SM75__
constexpr static bool allowable_aot_mode()
constexpr auto device_arch_compare_op_lt
static constexpr ext::oneapi::experimental::architecture min_amd_gpu_architecture
static constexpr std::optional< ext::oneapi::experimental::architecture > get_category_min_architecture(ext::oneapi::experimental::arch_category Category)
constexpr static bool device_architecture_comparison_aot(Compare comp)
constexpr static bool device_architecture_is_in_category_aot()
static constexpr ext::oneapi::experimental::architecture max_amd_gpu_architecture
static constexpr ext::oneapi::experimental::architecture min_intel_gpu_architecture
constexpr static bool device_architecture_is_in_categories()
static constexpr ext::oneapi::experimental::architecture min_nvidia_gpu_architecture
constexpr auto device_arch_compare_op_gt
constexpr static std::optional< ext::oneapi::experimental::arch_category > get_device_architecture_category(ext::oneapi::experimental::architecture arch)
constexpr static bool device_architecture_is()
constexpr static bool is_aot_for_architecture(ext::oneapi::experimental::architecture arch)
constexpr static std::optional< ext::oneapi::experimental::architecture > get_current_architecture_aot()
static constexpr std::optional< ext::oneapi::experimental::architecture > get_category_max_architecture(ext::oneapi::experimental::arch_category Category)
static constexpr ext::oneapi::experimental::architecture max_nvidia_gpu_architecture
constexpr auto device_arch_compare_op_ge
constexpr auto device_arch_compare_op_le
static constexpr bool is_allowable_aot_mode
static constexpr ext::oneapi::experimental::architecture max_intel_gpu_architecture
static constexpr int ConditionAnd
void call_if_on_device_conditionally(T fn, Condition...)
static constexpr int ConditionArchitecture
static constexpr int ConditionOr
static constexpr int ConditionNot
void call_if_on_device_conditionally_helper(T fn, std::integer_sequence< int, Is... >)
static constexpr int ConditionAspect
constexpr static auto if_architecture_is_lt(T fn)
The condition is true only if the device which executes the if_architecture_is_lt function has an arc...
constexpr static auto if_architecture_is(T fn)
The condition is true only if the device which executes the if_architecture_is function has any one o...
constexpr static auto if_architecture_is_between(T fn)
The condition is true only if the device which executes the if_architecture_is_between function has a...
constexpr static auto if_architecture_is_le(T fn)
The condition is true only if the device which executes the if_architecture_is_le function has an arc...
constexpr static auto if_architecture_is_ge(T fn)
The condition is true only if the device which executes the if_architecture_is_ge function has an arc...
constexpr static auto if_architecture_is_gt(T fn)
The condition is true only if the device which executes the if_architecture_is_gt function has an arc...
decltype(append< ConditionAnd >(rest1{}, rest2{})) seq
static auto append(std::integer_sequence< int, I1s... >, std::integer_sequence< int, I2s... >)
static auto append(std::integer_sequence< int, Is... >)
typename ConditionAnyArchitectureBuilder< Archs... >::seq rest
decltype(append< ConditionOr, ConditionArchitecture, arch >(rest{})) seq
std::integer_sequence< int, ConditionArchitecture, arch > seq
static auto append(std::integer_sequence< int, Is... >)
decltype(append< ConditionNot >(rest{})) seq
decltype(append< ConditionOr >(rest1{}, rest2{})) seq
static auto append(std::integer_sequence< int, I1s... >, std::integer_sequence< int, I2s... >)