Skip to content

[SYCL] Use uniform group instructions in headers #5705

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

Merged
merged 4 commits into from
Mar 22, 2022
Merged

[SYCL] Use uniform group instructions in headers #5705

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
9ed3ef2
[SYCL] Use uniform group instructions in headers
aobolensk Mar 1, 2022
f035628
Update instruction names
aobolensk Mar 3, 2022
121db8b
Address review comments
aobolensk Mar 3, 2022
8bd1e8c
Update check-clang test
aobolensk Mar 18, 2022
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
3 changes: 2 additions & 1 deletion clang/lib/Driver/ToolChains/Clang.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9188,7 +9188,8 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
ExtArg += ",+SPV_INTEL_token_type"
",+SPV_INTEL_bfloat16_conversion"
",+SPV_INTEL_joint_matrix"
",+SPV_INTEL_hw_thread_queries";
",+SPV_INTEL_hw_thread_queries"
",+SPV_KHR_uniform_group_instructions";
TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
}
for (auto I : Inputs) {
Expand Down
6 changes: 3 additions & 3 deletions clang/lib/Sema/SPIRVBuiltins.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -929,13 +929,13 @@ foreach name = ["GroupBroadcast"] in {
}
}

foreach name = ["GroupIAdd", "GroupNonUniformIMul", "GroupNonUniformBitwiseOr",
"GroupNonUniformBitwiseXor", "GroupNonUniformBitwiseAnd"] in {
foreach name = ["GroupIAdd", "GroupIMulKHR", "GroupBitwiseOrKHR",
"GroupBitwiseXorKHR", "GroupBitwiseAndKHR"] in {
def : SPVBuiltin<name, [AIGenTypeN, UInt, UInt, AIGenTypeN], Attr.Convergent>;
}

foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax",
"GroupNonUniformFMul"] in {
"GroupFMulKHR"] in {
def : SPVBuiltin<name, [FGenTypeN, UInt, UInt, FGenTypeN], Attr.Convergent>;
}

Expand Down
3 changes: 2 additions & 1 deletion clang/test/Driver/sycl-spirv-ext.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,8 @@
// CHECK-DEFAULT-SAME:,+SPV_INTEL_token_type
// CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
// CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
// CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries"
// CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries
// CHECK-DEFAULT-SAME:,+SPV_KHR_uniform_group_instructions"
// CHECK-FPGA-HW: llvm-spirv{{.*}}"-spirv-ext=-all
// CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_add
// CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_min_max
Expand Down
20 changes: 10 additions & 10 deletions sycl/include/sycl/ext/oneapi/functional.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,16 +82,16 @@ __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
__SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)

__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformIMul, sycl::multiplies<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformIMul, sycl::multiplies<T>)
__SYCL_CALC_OVERLOAD(GroupOpFP, NonUniformFMul, sycl::multiplies<T>)

__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseOr, sycl::bit_or<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseOr, sycl::bit_or<T>)
__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseXor, sycl::bit_xor<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseXor, sycl::bit_xor<T>)
__SYCL_CALC_OVERLOAD(GroupOpISigned, NonUniformBitwiseAnd, sycl::bit_and<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, NonUniformBitwiseAnd, sycl::bit_and<T>)
__SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies<T>)
__SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies<T>)

__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrKHR, sycl::bit_or<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrKHR, sycl::bit_or<T>)
__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorKHR, sycl::bit_xor<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorKHR, sycl::bit_xor<T>)
__SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndKHR, sycl::bit_and<T>)
__SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndKHR, sycl::bit_and<T>)

#undef __SYCL_CALC_OVERLOAD

Expand Down
Loading