Skip to content

Commit d4251e3

Browse files
AlexeySachkovAlexeySachkov
authored
[SYCL] Fix handling of multiple usages of composite spec constants (#2894)
Fixed the issue that instead of re-using previously assigned IDs for elements of a composite spec constant, all elements used the same ID, which was taken from the last element of the composite.
1 parent 74a68b7 commit d4251e3

File tree

3 files changed

+164
-20
lines changed

3 files changed

+164
-20
lines changed

‎llvm/test/tools/sycl-post-link/multiple-composite-spec-const-usages-2.ll

+78
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,78 @@
1+
; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
2+
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
3+
;
4+
; This test is intended to check that sycl-post-link tool is capable of handling
5+
; situations when the same composite specialization constants is used more than
6+
; once. Unlike multiple-composite-spec-const-usages.ll test, this is a real life
7+
; LLVM IR example
8+
;
9+
; CHECK-LABEL: @_ZTSN4test8kernel_tIfEE
10+
; CHECK: %[[#X1:]] = call float @_Z20__spirv_SpecConstantif(i32 0, float 0
11+
; CHECK: %[[#Y1:]] = call float @_Z20__spirv_SpecConstantif(i32 1, float 0
12+
; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X1]], float %[[#Y1]]), !SYCL_SPEC_CONST_SYM_ID ![[#ID:]]
13+
; CHECK-LABEL: @_ZTSN4test8kernel_tIiEE
14+
; CHECK: %[[#X2:]] = call float @_Z20__spirv_SpecConstantif(i32 0, float 0
15+
; CHECK: %[[#Y2:]] = call float @_Z20__spirv_SpecConstantif(i32 1, float 0
16+
; CHECK: call {{.*}} @_Z29__spirv_SpecConstantCompositeff(float %[[#X2]], float %[[#Y2]]), !SYCL_SPEC_CONST_SYM_ID ![[#ID]]
17+
; CHECK: ![[#ID]] = !{!"_ZTS11sc_kernel_t", i32 0, i32 1}
18+
19+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
20+
target triple = "spir64-unknown-unknown-sycldevice"
21+
22+
%"struct._ZTSN4test5pod_tE.test::pod_t" = type { float, float }
23+
24+
$_ZTSN4test8kernel_tIfEE = comdat any
25+
26+
$_ZTSN4test8kernel_tIiEE = comdat any
27+
28+
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv = private unnamed_addr addrspace(1) constant [18 x i8] c"_ZTS11sc_kernel_t\00", align 1
29+
30+
; Function Attrs: convergent norecurse
31+
define weak_odr dso_local spir_kernel void @_ZTSN4test8kernel_tIfEE() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
32+
entry:
33+
%ref.tmp.i = alloca %"struct._ZTSN4test5pod_tE.test::pod_t", align 4
34+
%0 = bitcast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to i8*
35+
call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0) #3
36+
%1 = addrspacecast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to %"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)*
37+
call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4
38+
call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0) #3
39+
ret void
40+
}
41+
42+
; Function Attrs: argmemonly nofree nosync nounwind willreturn
43+
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
44+
45+
; Function Attrs: argmemonly nofree nosync nounwind willreturn
46+
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
47+
48+
; Function Attrs: convergent
49+
declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4, i8 addrspace(4)*) local_unnamed_addr #2
50+
51+
; Function Attrs: convergent norecurse
52+
define weak_odr dso_local spir_kernel void @_ZTSN4test8kernel_tIiEE() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
53+
entry:
54+
%ref.tmp.i = alloca %"struct._ZTSN4test5pod_tE.test::pod_t", align 4
55+
%0 = bitcast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to i8*
56+
call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0) #3
57+
%1 = addrspacecast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to %"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)*
58+
call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4
59+
call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0) #3
60+
ret void
61+
}
62+
63+
attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="repro-1.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
64+
attributes #1 = { argmemonly nofree nosync nounwind willreturn }
65+
attributes #2 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
66+
attributes #3 = { nounwind }
67+
attributes #4 = { convergent }
68+
69+
!llvm.module.flags = !{!0}
70+
!opencl.spir.version = !{!1}
71+
!spirv.Source = !{!2}
72+
!llvm.ident = !{!3}
73+
74+
!0 = !{i32 1, !"wchar_size", i32 4}
75+
!1 = !{i32 1, i32 2}
76+
!2 = !{i32 4, i32 100000}
77+
!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 9b7086f7cef079b80ac5e137394f8d77d5d49c3e)"}
78+
!4 = !{}

‎llvm/tools/sycl-post-link/SpecConstants.cpp

+10-20
Original file line numberDiff line numberDiff line change
@@ -318,12 +318,7 @@ Instruction *emitSpecConstantComposite(Type *Ty,
318318
/// first ID. If \c IsNewSpecConstant is false, this vector is expected to
319319
/// contain enough elements to assign ID to each scalar element encountered in
320320
/// the specified composite type.
321-
/// @param IsNewSpecConstant [in] Flag to specify whether \c IDs vector should
322-
/// be filled with new IDs or it should be used as-is to replicate an existing
323-
/// spec constant
324-
/// @param [in,out] IsFirstElement Flag indicating whether this function is
325-
/// handling the first scalar element encountered in the specified composite
326-
/// type \c Ty or not.
321+
/// @param [in,out] Index Index of scalar element within a composite type
327322
///
328323
/// @returns Instruction* representing specialization constant in LLVM IR, which
329324
/// is in SPIR-V friendly LLVM IR form.
@@ -335,22 +330,20 @@ Instruction *emitSpecConstantComposite(Type *Ty,
335330
/// encountered scalars and assigns them IDs (or re-uses existing ones).
336331
Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore,
337332
SmallVectorImpl<unsigned> &IDs,
338-
bool IsNewSpecConstant,
339-
bool &IsFirstElement) {
333+
unsigned &Index) {
340334
if (!Ty->isArrayTy() && !Ty->isStructTy() && !Ty->isVectorTy()) { // Scalar
341-
if (IsNewSpecConstant && !IsFirstElement) {
335+
if (Index >= IDs.size()) {
342336
// If it is a new specialization constant, we need to generate IDs for
343337
// scalar elements, starting with the second one.
344338
IDs.push_back(IDs.back() + 1);
345339
}
346-
IsFirstElement = false;
347-
return emitSpecConstant(IDs.back(), Ty, InsertBefore);
340+
return emitSpecConstant(IDs[Index++], Ty, InsertBefore);
348341
}
349342

350343
SmallVector<Instruction *, 8> Elements;
351344
auto LoopIteration = [&](Type *Ty) {
352-
Elements.push_back(emitSpecConstantRecursiveImpl(
353-
Ty, InsertBefore, IDs, IsNewSpecConstant, IsFirstElement));
345+
Elements.push_back(
346+
emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index));
354347
};
355348

356349
if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
@@ -374,11 +367,9 @@ Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore,
374367

375368
/// Wrapper intended to hide IsFirstElement argument from the caller
376369
Instruction *emitSpecConstantRecursive(Type *Ty, Instruction *InsertBefore,
377-
SmallVectorImpl<unsigned> &IDs,
378-
bool IsNewSpecConstant) {
379-
bool IsFirstElement = true;
380-
return emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, IsNewSpecConstant,
381-
IsFirstElement);
370+
SmallVectorImpl<unsigned> &IDs) {
371+
unsigned Index = 0;
372+
return emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index);
382373
}
383374

384375
} // namespace
@@ -446,8 +437,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
446437

447438
// 3. Transform to spirv intrinsic _Z*__spirv_SpecConstant* or
448439
// _Z*__spirv_SpecConstantComposite
449-
auto *SPIRVCall =
450-
emitSpecConstantRecursive(SCTy, CI, IDs, IsNewSpecConstant);
440+
auto *SPIRVCall = emitSpecConstantRecursive(SCTy, CI, IDs);
451441
if (IsNewSpecConstant) {
452442
// emitSpecConstantRecursive might emit more than one spec constant
453443
// (because of composite types) and therefore, we need to ajudst

‎sycl/test/on-device/spec_const/multiple-usages-of-composite.cpp

+76
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,76 @@
1+
// UNSUPPORTED: cuda
2+
//
3+
// RUN: %clangxx -fsycl %s -o %t.out
4+
// RUN: %RUN_ON_HOST %t.out | FileCheck %s
5+
// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
6+
// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
7+
//
8+
// The test checks that multiple usages of the same specialization constant
9+
// works correctly: toolchain processes them correctly and runtime can
10+
// correctly execute the program.
11+
//
12+
// CHECK: --------> 1
13+
14+
#include <CL/sycl.hpp>
15+
16+
using namespace cl::sycl;
17+
18+
class sc_kernel_t;
19+
20+
namespace test {
21+
22+
struct pod_t {
23+
float x;
24+
float y;
25+
};
26+
27+
template <typename T> class kernel_t {
28+
public:
29+
using sc_t = sycl::ONEAPI::experimental::spec_constant<pod_t, sc_kernel_t>;
30+
31+
kernel_t(const sc_t &sc, cl::sycl::stream &strm) : sc_(sc), strm_(strm) {}
32+
33+
void operator()(cl::sycl::id<1> i) const {
34+
strm_ << "--------> " << sc_.get().x << sycl::endl;
35+
}
36+
37+
sc_t sc_;
38+
cl::sycl::stream strm_;
39+
};
40+
41+
template <typename T> class kernel_driver_t {
42+
public:
43+
void execute(const pod_t &pod) {
44+
device dev = sycl::device(default_selector{});
45+
context ctx = context(dev);
46+
queue q(dev);
47+
48+
cl::sycl::program p(q.get_context());
49+
auto sc = p.set_spec_constant<sc_kernel_t>(pod);
50+
p.build_with_kernel_type<kernel_t<T>>();
51+
52+
q.submit([&](cl::sycl::handler &cgh) {
53+
cl::sycl::stream strm(1024, 256, cgh);
54+
kernel_t<T> func(sc, strm);
55+
56+
auto sycl_kernel = p.get_kernel<kernel_t<T>>();
57+
cgh.parallel_for(sycl_kernel, cl::sycl::range<1>(1), func);
58+
});
59+
q.wait();
60+
}
61+
};
62+
63+
template class kernel_driver_t<float>;
64+
65+
// The line below instantiates the second use of the spec constant named
66+
// `sc_kernel_t`, which used to corrupt the spec constant content
67+
template class kernel_driver_t<int>;
68+
} // namespace test
69+
70+
int main() {
71+
test::pod_t pod = {1, 2};
72+
test::kernel_driver_t<float> kd_float;
73+
kd_float.execute(pod);
74+
75+
return 0;
76+
}

0 commit comments

Comments
 (0)