Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[DOC][ABI-BREAK] Remove bfloat16 math aspect. #13351

Merged
merged 10 commits into from
Jul 3, 2024
Merged

[DOC][ABI-BREAK] Remove bfloat16 math aspect. #13351

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
33e7a6c
Remove ref to bfloat16 math aspect.
JackAKirk Apr 10, 2024
348f2c8
Remove bfloat16_math pt1.
JackAKirk Apr 11, 2024
c7d2e74
Fix format.
JackAKirk Apr 11, 2024
3cf0191
Check that conversions give same error.
JackAKirk Apr 12, 2024
be3ba04
Merge branch 'sycl' into bfloat16-math-remove-aspect
JackAKirk Jul 1, 2024
8afcaa8
Remove aspect from deviceConfigfile.
JackAKirk Jul 2, 2024
4f27e0b
Don't change spirv tests.
JackAKirk Jul 2, 2024
81d4401
Merge branch 'sycl' into bfloat16-math-remove-aspect
JackAKirk Jul 2, 2024
791d68d
Merge branch 'sycl' into bfloat16-math-remove-aspect
JackAKirk Jul 2, 2024
b723804
Merge branch 'sycl' into bfloat16-math-remove-aspect
JackAKirk Jul 3, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 7 additions & 8 deletions llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,6 @@ def AspectExt_oneapi_native_assert : Aspect<"ext_oneapi_native_assert">;
def AspectHost_debuggable : Aspect<"host_debuggable">;
def AspectExt_intel_gpu_hw_threads_per_eu : Aspect<"ext_intel_gpu_hw_threads_per_eu">;
def AspectExt_oneapi_cuda_async_barrier : Aspect<"ext_oneapi_cuda_async_barrier">;
def AspectExt_oneapi_bfloat16_math_functions : Aspect<"ext_oneapi_bfloat16_math_functions">;
def AspectExt_intel_free_memory : Aspect<"ext_intel_free_memory">;
def AspectExt_intel_device_id : Aspect<"ext_intel_device_id">;
def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate">;
Expand Down Expand Up @@ -125,7 +124,7 @@ def : TargetInfo<"__TestAspectList",
AspectExt_intel_max_mem_bandwidth, AspectExt_intel_mem_channel, AspectUsm_atomic_host_allocations,
AspectUsm_atomic_shared_allocations, AspectAtomic64, AspectExt_intel_device_info_uuid, AspectExt_oneapi_srgb,
AspectExt_oneapi_native_assert, AspectHost_debuggable, AspectExt_intel_gpu_hw_threads_per_eu,
AspectExt_oneapi_cuda_async_barrier, AspectExt_oneapi_bfloat16_math_functions, AspectExt_intel_free_memory,
AspectExt_oneapi_cuda_async_barrier, AspectExt_intel_free_memory,
AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectEmulated,
AspectExt_intel_legacy_image, AspectExt_oneapi_bindless_images,
AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm,
Expand Down Expand Up @@ -198,17 +197,17 @@ def : CudaTargetInfo<"nvidia_gpu_sm_70", !listconcat(CudaMinAspects, CudaBindles
def : CudaTargetInfo<"nvidia_gpu_sm_72", !listconcat(CudaMinAspects, CudaBindlessImagesAspects, [AspectFp16, AspectAtomic64])>;
def : CudaTargetInfo<"nvidia_gpu_sm_75", !listconcat(CudaMinAspects, CudaBindlessImagesAspects, [AspectFp16, AspectAtomic64])>;
def : CudaTargetInfo<"nvidia_gpu_sm_80", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;
def : CudaTargetInfo<"nvidia_gpu_sm_86", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;
def : CudaTargetInfo<"nvidia_gpu_sm_87", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;
def : CudaTargetInfo<"nvidia_gpu_sm_89", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;
def : CudaTargetInfo<"nvidia_gpu_sm_90", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;
def : CudaTargetInfo<"nvidia_gpu_sm_90a", !listconcat(CudaMinAspects, CudaBindlessImagesAspects,
[AspectFp16, AspectAtomic64, AspectExt_oneapi_bfloat16_math_functions, AspectExt_oneapi_cuda_async_barrier])>;
[AspectFp16, AspectAtomic64, AspectExt_oneapi_cuda_async_barrier])>;

//
// HIP / AMDGPU device aspects
Expand Down
1 change: 0 additions & 1 deletion sycl/doc/design/DeviceConfigFile.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -176,7 +176,6 @@ def AspectExt_oneapi_native_assert : Aspect<"ext_oneapi_native_assert">;
def AspectHost_debuggable : Aspect<"host_debuggable">;
def AspectExt_intel_gpu_hw_threads_per_eu : Aspect<"ext_intel_gpu_hw_threads_per_eu">;
def AspectExt_oneapi_cuda_async_barrier : Aspect<"ext_oneapi_cuda_async_barrier">;
def AspectExt_oneapi_bfloat16_math_functions : Aspect<"ext_oneapi_bfloat16_math_functions">;
def AspectExt_intel_free_memory : Aspect<"ext_intel_free_memory">;
def AspectExt_intel_device_id : Aspect<"ext_intel_device_id">;
def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate">;
Expand Down
21 changes: 6 additions & 15 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,6 +67,12 @@ The descriptions of the `fma`, `fmin`, `fmax`, `fabs`, `isnan`, `ceil`, `floor`,
specification:
https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_math_functions.

[NOTE]
The bfloat16 type is supported on all devices. DPC++ currently supports some
bfloat16 type math functions natively on Intel Xe HP GPUs and Nvidia GPUs with
Compute Capability >= SM80. On other devices, and in host code, such functions
are emulated in software.

== Specification

=== Feature test macro
Expand All @@ -86,21 +92,6 @@ supports.
|1 |The APIs of this experimental extension are not versioned, so the feature-test macro always has this value.
|===

=== Extension to `enum class aspect`

[source]
----
namespace sycl {
enum class aspect {
...
sycl_ext_oneapi_bfloat16_math_functions
}
}
----

If a SYCL device has the `sycl_ext_oneapi_bfloat16_math_functions` aspect,
then it supports the `bfloat16` math functions described in the next section.

=== Math Functions

==== isnan
Expand Down
10 changes: 0 additions & 10 deletions sycl/include/sycl/device_aspect_macros.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -183,11 +183,6 @@
#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_cuda_async_barrier__ 0
#endif

#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bfloat16_math_functions__
// __SYCL_ASPECT(ext_oneapi_bfloat16_math_functions, 35)
#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bfloat16_math_functions__ 0
#endif

#ifndef __SYCL_ALL_DEVICES_HAVE_ext_intel_free_memory__
// __SYCL_ASPECT(ext_intel_free_memory, 36)
#define __SYCL_ALL_DEVICES_HAVE_ext_intel_free_memory__ 0
Expand Down Expand Up @@ -561,11 +556,6 @@
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_cuda_async_barrier__ 0
#endif

#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bfloat16_math_functions__
// __SYCL_ASPECT(ext_oneapi_bfloat16_math_functions, 35)
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bfloat16_math_functions__ 0
#endif

#ifndef __SYCL_ANY_DEVICE_HAS_ext_intel_free_memory__
// __SYCL_ASPECT(ext_intel_free_memory, 36)
#define __SYCL_ANY_DEVICE_HAS_ext_intel_free_memory__ 0
Expand Down
1 change: 0 additions & 1 deletion sycl/include/sycl/info/aspects.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,6 @@ __SYCL_ASPECT(ext_oneapi_native_assert, 31)
__SYCL_ASPECT(host_debuggable, 32)
__SYCL_ASPECT(ext_intel_gpu_hw_threads_per_eu, 33)
__SYCL_ASPECT(ext_oneapi_cuda_async_barrier, 34)
__SYCL_ASPECT(ext_oneapi_bfloat16_math_functions, 35)
__SYCL_ASPECT(ext_intel_free_memory, 36)
__SYCL_ASPECT(ext_intel_device_id, 37)
__SYCL_ASPECT(ext_intel_memory_clock_rate, 38)
Expand Down
2 changes: 0 additions & 2 deletions sycl/include/sycl/info/device_traits.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,8 +202,6 @@ __SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_srgb, bool,
PI_DEVICE_INFO_IMAGE_SRGB)
__SYCL_PARAM_TRAITS_SPEC(device, ext_intel_mem_channel, bool,
PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT)
__SYCL_PARAM_TRAITS_SPEC(device, ext_oneapi_bfloat16_math_functions, bool,
PI_EXT_ONEAPI_DEVICE_INFO_BFLOAT16_MATH_FUNCTIONS)

//Deprecated oneapi/intel extension
//TODO:Remove when possible
Expand Down
2 changes: 0 additions & 2 deletions sycl/source/detail/device_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,8 +346,6 @@ bool device_impl::has(aspect Aspect) const {
return has_extension("cl_khr_fp16");
case aspect::fp64:
return has_extension("cl_khr_fp64");
case aspect::ext_oneapi_bfloat16_math_functions:
return get_info<info::device::ext_oneapi_bfloat16_math_functions>();
case aspect::int64_base_atomics:
return has_extension("cl_khr_int64_base_atomics");
case aspect::int64_extended_atomics:
Expand Down
19 changes: 0 additions & 19 deletions sycl/source/detail/device_info.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -301,25 +301,6 @@ struct get_device_info_impl<std::vector<memory_scope>,
}
};

// Specialization for bf16 math functions
template <>
struct get_device_info_impl<bool,
info::device::ext_oneapi_bfloat16_math_functions> {
static bool get(const DeviceImplPtr &Dev) {
bool result = false;

sycl::detail::pi::PiResult Err =
Dev->getPlugin()->call_nocheck<PiApiKind::piDeviceGetInfo>(
Dev->getHandleRef(),
PiInfoCode<info::device::ext_oneapi_bfloat16_math_functions>::value,
sizeof(result), &result, nullptr);
if (Err != PI_SUCCESS) {
return false;
}
return result;
}
};

// Specialization for exec_capabilities, OpenCL returns a bitfield
template <>
struct get_device_info_impl<std::vector<info::execution_capability>,
Expand Down
1 change: 0 additions & 1 deletion sycl/test-e2e/BFloat16/bfloat16_conversions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,6 @@
// UNSUPPORTED: accelerator

// FIXME: enable opaque pointers support on CPU.
// UNSUPPORTED: cpu

//==---------- bfloat16_conversions.cpp - SYCL bfloat16 type test ---------==//
//
Expand Down
1 change: 0 additions & 1 deletion sycl/test-e2e/BFloat16/bfloat16_type.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,6 @@
// UNSUPPORTED: accelerator

// FIXME: enable opaque pointers support on CPU.
// UNSUPPORTED: cpu

//==----------- bfloat16_type.cpp - SYCL bfloat16 type test ----------------==//
//
Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/Basic/aspects.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,9 +51,6 @@ int main() {
if (plt.has(aspect::fp64)) {
std::cout << " fp64" << std::endl;
}
if (plt.has(aspect::ext_oneapi_bfloat16_math_functions)) {
std::cout << " ext_oneapi_bfloat16_math_functions" << std::endl;
}
if (plt.has(aspect::atomic64)) {
std::cout << " atomic64" << std::endl;
}
Expand Down
1 change: 0 additions & 1 deletion sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3983,7 +3983,6 @@ _ZNK4sycl3_V16device13get_info_implINS0_4info6device32atomic_memory_scope_capabi
_ZNK4sycl3_V16device13get_info_implINS0_4info6device33ext_intel_gpu_subslices_per_sliceEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device33ext_oneapi_max_global_work_groupsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device33usm_restricted_shared_allocationsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device34ext_oneapi_bfloat16_math_functionsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device35ext_intel_gpu_eu_count_per_subsliceEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device38sub_group_independent_forward_progressEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
_ZNK4sycl3_V16device13get_info_implINS0_4info6device4nameEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv
Expand Down
1 change: 0 additions & 1 deletion sycl/test/abi/sycl_symbols_windows.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -137,7 +137,6 @@
??$get_info_impl@Uext_intel_max_mem_bandwidth@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ
??$get_info_impl@Uext_intel_mem_channel@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_NXZ
??$get_info_impl@Uext_intel_pci_address@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AVstring@detail@12@XZ
??$get_info_impl@Uext_oneapi_bfloat16_math_functions@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_NXZ
??$get_info_impl@Uext_oneapi_max_global_work_groups@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ
??$get_info_impl@Uext_oneapi_max_work_groups_1d@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$id@$00@12@XZ
??$get_info_impl@Uext_oneapi_max_work_groups_2d@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$id@$01@12@XZ
Expand Down
8 changes: 3 additions & 5 deletions sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,8 +8,7 @@ using namespace sycl;
using namespace ext::oneapi::experimental;

static constexpr auto device_has_all = device_has<
aspect::ext_oneapi_cuda_async_barrier,
aspect::ext_oneapi_bfloat16_math_functions, aspect::custom, aspect::fp16,
aspect::ext_oneapi_cuda_async_barrier, aspect::custom, aspect::fp16,
aspect::fp64, aspect::image, aspect::online_compiler, aspect::online_linker,
aspect::queue_profiling, aspect::usm_device_allocations,
aspect::usm_system_allocations, aspect::ext_intel_pci_address, aspect::cpu,
Expand Down Expand Up @@ -131,7 +130,6 @@ int main() {
}

// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_cuda_async_barrier", i32 [[ext_oneapi_cuda_async_barrier_ASPECT_MD:[0-9]+]]}
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_bfloat16_math_functions", i32 [[ext_oneapi_bfloat16_math_functions_ASPECT_MD:[0-9]+]]}
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"custom", i32 [[custom_ASPECT_MD:[0-9]+]]}
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp16", i32 [[fp16_ASPECT_MD:[0-9]+]]}
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp64", i32 [[fp64_ASPECT_MD:[0-9]+]]}
Expand Down Expand Up @@ -165,5 +163,5 @@ int main() {
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_free_memory", i32 [[ext_intel_free_memory_ASPECT_MD:[0-9]+]]}
// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_device_id", i32 [[ext_intel_device_id_ASPECT_MD:[0-9]+]]}

// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
// CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
// CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
Loading
Loading