Skip to content

Commit e9c8acf

Browse files
committed
[SYCL] Fix crash in clang for SYCL kernels with accessors and SC
In case a SYCL kernel captures specialization constants and more then one accessor it crashes. I think the crash happens due to an extra increment of KernelFuncParam, in CreateOpenCLKernelBody() method. Signed-off-by: Alexey Sotkin <alexey.sotkin@intel.com>
1 parent 16fffef commit e9c8acf

File tree

2 files changed

+47
-0
lines changed

2 files changed

+47
-0
lines changed

clang/lib/Sema/SemaSYCL.cpp

+1
Original file line numberDiff line numberDiff line change
@@ -815,6 +815,7 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
815815
InitMethodName, BodyStmts);
816816
} else if (Util::isSyclSpecConstantType(FieldType)) {
817817
// Just skip specialization constants - not part of signature.
818+
continue;
818819
} else if (CRD || FieldType->isScalarType()) {
819820
// If field has built-in or a structure/class type just initialize
820821
// this field with corresponding kernel argument using copy
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: %clangxx -fsycl-device-only %s -o %t.bc
2+
//
3+
//==----------- spec_const_accessors.cpp -----------------------------------==//
4+
//
5+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
6+
// See https://llvm.org/LICENSE.txt for license information.
7+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
8+
//
9+
//===----------------------------------------------------------------------===//
10+
// The test checks that Clang doesn't crash if multiple accessors are used
11+
// together with specialization constants in single kernel.
12+
13+
#include <CL/sycl.hpp>
14+
15+
using namespace sycl;
16+
17+
class MyInt32Const;
18+
class MyKernel;
19+
20+
int32_t val = 10;
21+
22+
int main() {
23+
cl::sycl::queue queue;
24+
cl::sycl::program program(queue.get_context());
25+
26+
cl::sycl::experimental::spec_constant<int32_t, MyInt32Const> i32 =
27+
program.set_spec_constant<MyInt32Const>(val);
28+
29+
program.build_with_kernel_type<MyKernel>();
30+
int32_t a = 0;
31+
int32_t b = 0;
32+
cl::sycl::buffer<int32_t, 1> buf_a(&a, 1);
33+
cl::sycl::buffer<int32_t, 1> buf_b(&b, 1);
34+
35+
queue.submit([&](cl::sycl::handler &cgh) {
36+
auto acc_a = buf_a.get_access<cl::sycl::access::mode::write>(cgh);
37+
auto acc_b = buf_b.get_access<cl::sycl::access::mode::write>(cgh);
38+
cgh.single_task<MyKernel>(
39+
program.get_kernel<MyKernel>(),
40+
[=]() {
41+
acc_a[0] = i32.get();
42+
acc_b[0] = i32.get();
43+
});
44+
});
45+
return 0;
46+
}

0 commit comments

Comments
 (0)