diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 54357d1377c77..bbaad03a27479 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -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">; @@ -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, @@ -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 diff --git a/sycl/doc/design/DeviceConfigFile.md b/sycl/doc/design/DeviceConfigFile.md index ba9c0cebbfe4a..505c65226a809 100644 --- a/sycl/doc/design/DeviceConfigFile.md +++ b/sycl/doc/design/DeviceConfigFile.md @@ -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">; diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc index 6359515a67b9d..11dade8c226f0 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc @@ -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 @@ -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 diff --git a/sycl/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp index d756b0a62e88a..10f61c1e48435 100644 --- a/sycl/include/sycl/device_aspect_macros.hpp +++ b/sycl/include/sycl/device_aspect_macros.hpp @@ -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 @@ -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 diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 3b744a89dbb90..2151776ddb8e7 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -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) diff --git a/sycl/include/sycl/info/device_traits.def b/sycl/include/sycl/info/device_traits.def index dc50c5e920502..730ef88e335d4 100644 --- a/sycl/include/sycl/info/device_traits.def +++ b/sycl/include/sycl/info/device_traits.def @@ -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 diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 8d39a61027208..e22d87be839ae 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -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(); case aspect::int64_base_atomics: return has_extension("cl_khr_int64_base_atomics"); case aspect::int64_extended_atomics: diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index 4a1a410c6845d..98d9d171f5737 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -301,25 +301,6 @@ struct get_device_info_impl, } }; -// Specialization for bf16 math functions -template <> -struct get_device_info_impl { - static bool get(const DeviceImplPtr &Dev) { - bool result = false; - - sycl::detail::pi::PiResult Err = - Dev->getPlugin()->call_nocheck( - Dev->getHandleRef(), - PiInfoCode::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, diff --git a/sycl/test-e2e/BFloat16/bfloat16_conversions.cpp b/sycl/test-e2e/BFloat16/bfloat16_conversions.cpp index cb59576a2eeb7..d6fa57648ca6c 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_conversions.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_conversions.cpp @@ -6,7 +6,6 @@ // UNSUPPORTED: accelerator // FIXME: enable opaque pointers support on CPU. -// UNSUPPORTED: cpu //==---------- bfloat16_conversions.cpp - SYCL bfloat16 type test ---------==// // diff --git a/sycl/test-e2e/BFloat16/bfloat16_type.cpp b/sycl/test-e2e/BFloat16/bfloat16_type.cpp index ccd04b5c54188..73936eb5dc4e7 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_type.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_type.cpp @@ -7,7 +7,6 @@ // UNSUPPORTED: accelerator // FIXME: enable opaque pointers support on CPU. -// UNSUPPORTED: cpu //==----------- bfloat16_type.cpp - SYCL bfloat16 type test ----------------==// // diff --git a/sycl/test-e2e/Basic/aspects.cpp b/sycl/test-e2e/Basic/aspects.cpp index 75daf263e638e..bd47d1b4bf1cb 100644 --- a/sycl/test-e2e/Basic/aspects.cpp +++ b/sycl/test-e2e/Basic/aspects.cpp @@ -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; } diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 4a7d0dcdcb0eb..c8d713a57c1e7 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -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 diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 00d45e8913778..adc189f5eb5e7 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -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 diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index aa1988837d8c4..055b25b920b8b 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -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, @@ -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]+]]} @@ -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]]" diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has_macro.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has_macro.cpp index 189d29184d3c5..e4ff52bf02b99 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has_macro.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has_macro.cpp @@ -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, @@ -37,7 +36,6 @@ SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( (device_has)) void Func2() {} // 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]+]]} @@ -73,6 +71,6 @@ SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( // CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_memory_clock_rate", i32 [[ext_intel_memory_clock_rate_ASPECT_MD:[0-9]+]]} // CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_memory_bus_width", i32 [[ext_intel_memory_bus_width_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]],[[ext_intel_memory_clock_rate_ASPECT_MD]],[[ext_intel_memory_bus_width_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]],[[ext_intel_memory_clock_rate_ASPECT_MD]],[[ext_intel_memory_bus_width_ASPECT_MD]]" // CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has" {{.*}} // CHECK-IR-DAG: attributes #[[DHAttr3]] = { {{.*}}"sycl-device-has"="[[fp16_ASPECT_MD]],[[atomic64_ASPECT_MD]]" diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp index 65717c2200628..3868c23f7535c 100644 --- a/sycl/test/extensions/properties/properties_kernel.cpp +++ b/sycl/test/extensions/properties/properties_kernel.cpp @@ -24,7 +24,6 @@ using device_has_all = aspect::ext_oneapi_native_assert, aspect::host_debuggable, aspect::ext_intel_gpu_hw_threads_per_eu, aspect::ext_oneapi_cuda_async_barrier, - aspect::ext_oneapi_bfloat16_math_functions, aspect::ext_intel_free_memory, aspect::ext_intel_device_id, aspect::ext_intel_memory_clock_rate, aspect::ext_intel_memory_bus_width>); @@ -116,7 +115,6 @@ int main() { singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); - singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); @@ -128,7 +126,7 @@ int main() { static_assert(is_property_value::value); static_assert(std::is_same_v); - static_assert(device_has_all::value.size() == 36); + static_assert(device_has_all::value.size() == 35); static_assert(device_has_all::value[0] == aspect::cpu); static_assert(device_has_all::value[1] == aspect::gpu); static_assert(device_has_all::value[2] == aspect::accelerator); @@ -169,13 +167,11 @@ int main() { aspect::ext_intel_gpu_hw_threads_per_eu); static_assert(device_has_all::value[30] == aspect::ext_oneapi_cuda_async_barrier); - static_assert(device_has_all::value[31] == - aspect::ext_oneapi_bfloat16_math_functions); - static_assert(device_has_all::value[32] == aspect::ext_intel_free_memory); - static_assert(device_has_all::value[33] == aspect::ext_intel_device_id); - static_assert(device_has_all::value[34] == + static_assert(device_has_all::value[31] == aspect::ext_intel_free_memory); + static_assert(device_has_all::value[32] == aspect::ext_intel_device_id); + static_assert(device_has_all::value[33] == aspect::ext_intel_memory_clock_rate); - static_assert(device_has_all::value[35] == + static_assert(device_has_all::value[34] == aspect::ext_intel_memory_bus_width); return 0; }