Skip to content

Commit

Permalink
[DOC][ABI-BREAK] Remove bfloat16 math aspect. (#13351)
Browse files Browse the repository at this point in the history 
This aspect is not required because bfloat16 math functions are
implemented for all devices via generic implementations. This PR updates
this status inline with the main bfloat16 extension/doc.

---------

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
  • Loading branch information
@JackAKirk
JackAKirk committed Jul 3, 2024
1 parent cb2efef commit f958dce
Show file tree
Hide file tree
Showing 16 changed files with 23 additions and 83 deletions.
15 changes: 7 additions & 8 deletions llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
View file
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
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
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
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
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
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
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
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
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
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
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
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
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
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

0 comments on commit f958dce

Please sign in to comment.