Skip to content

WIP Replace local clone with SYCL kernel object #1568

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

Closed
wants to merge 6 commits into from
Closed

WIP Replace local clone with SYCL kernel object #1568

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
b2f3950
[SYCL] Add kernel object to Kernel function and integration header
elizabethandrews Apr 22, 2020
3eae199
Fix lit failures
elizabethandrews Apr 27, 2020
5c643d6
Replace local clone Part 1
elizabethandrews Apr 27, 2020
cce4455
Clang format changes
elizabethandrews Apr 29, 2020
11aa9a2
Resolved merge conflicts
elizabethandrews Apr 29, 2020
4b8a799
Fix build error
elizabethandrews Apr 29, 2020
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
162 changes: 90 additions & 72 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -628,6 +628,16 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) {
Ctx.getTrivialTypeSourceInfo(Ty));
}

// Creates a parameter descriptor for kernel object
static ParamDesc makeParamDesc(const CXXRecordDecl *Src, QualType Ty) {
ASTContext &Ctx = Src->getASTContext();
// There is no name available for lambda object. Name for all
// kernel types (lambda and functor) is set as _arg_kernelObject.
std::string Name = "_arg_kernelObject";
return std::make_tuple(Ty, &Ctx.Idents.get(Name),
Ctx.getTrivialTypeSourceInfo(Ty));
}

static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src,
QualType Ty) {
// TODO: There is no name for the base available, but duplicate names are
Expand Down Expand Up @@ -721,12 +731,17 @@ static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent,
// A visitor function that dispatches to functions as defined in
// SyclKernelFieldHandler for the purposes of kernel generation.
template <typename... Handlers>
static void VisitRecordFields(RecordDecl::field_range Fields,
static void VisitRecordFields(CXXRecordDecl *KernelObject,
Handlers &... handlers) {

QualType KernelType = QualType(KernelObject->getTypeForDecl(), 0);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Are these only allowed to be top-level things? Can you clarify how this feature is supposed to work? I don't get it yet.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We're adding the kernel object as a parameter of the kernel function (just like we did with the fields). Currently in the kernel body, a kernel object clone is generated and initialized using the fields of kernel function. The clone is then subsequently used in kernel body code. We're now exploring the possibility of getting rid of the clone entirely and using this kernel object

(void)std::initializer_list<int>{
(handlers.handleKernelObject(KernelObject, KernelType), 0)...};

#define KF_FOR_EACH(FUNC) \
(void)std::initializer_list<int> { (handlers.FUNC(Field, FieldTy), 0)... }

for (const auto &Field : Fields) {
for (const auto &Field : KernelObject->fields()) {
QualType FieldTy = Field->getType();

if (Util::isSyclAccessorType(FieldTy))
Expand Down Expand Up @@ -781,6 +796,7 @@ template <typename Derived> class SyclKernelFieldHandler {
virtual void handlePointerType(FieldDecl *, QualType) {}
virtual void handleArrayType(FieldDecl *, QualType) {}
virtual void handleScalarType(FieldDecl *, QualType) {}
virtual void handleKernelObject(CXXRecordDecl *, QualType) {}
// Most handlers shouldn't be handling this, just the field checker.
virtual void handleOtherType(FieldDecl *, QualType) {}

Expand Down Expand Up @@ -830,6 +846,16 @@ class SyclKernelFieldChecker
<< 1 << FieldTy;
}
}
void handleKernelObject(CXXRecordDecl *KernelObject,
QualType KernelType) final {
// TODO: Is this check correct? SYCL spec only talks about kernel defined as
// named function objects. What about lambda functions?
/*if (!KernelObject->isStandardLayoutType())
IsInvalid =
Diag.Report(KernelObject->getLocation(),
diag::err_sycl_non_std_layout_type)
<< KernelType;*/
}

// We should be able to handle this, so we made it part of the visitor, but
// this is 'to be implemented'.
Expand Down Expand Up @@ -860,6 +886,11 @@ class SyclKernelDeclCreator
addParam(newParamDesc, FieldTy);
}

void addParam(const CXXRecordDecl *KernelObject, QualType KernelType) {
ParamDesc newParamDesc = makeParamDesc(KernelObject, KernelType);
addParam(newParamDesc, KernelType);
}

void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) {
ParamDesc newParamDesc =
makeParamDesc(SemaRef.getASTContext(), BS, FieldTy);
Expand Down Expand Up @@ -985,12 +1016,9 @@ class SyclKernelDeclCreator
addParam(FD, ModTy);
}

void handleScalarType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy);
}

// TODO: Accessors in structs
void handleStructType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy);
// addParam(FD, FieldTy);
}

void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
Expand All @@ -1002,6 +1030,11 @@ class SyclKernelDeclCreator
// See https://github.com/intel/llvm/issues/1552
}

void handleKernelObject(CXXRecordDecl *KernelObject,
QualType KernelType) final {
addParam(KernelObject, KernelType);
}

void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); }

FunctionDecl *getKernelDecl() { return KernelDecl; }
Expand All @@ -1010,16 +1043,15 @@ class SyclKernelDeclCreator
return ArrayRef<ParmVarDecl *>(std::begin(Params) + LastParamIndex,
std::end(Params));
}

ParmVarDecl *getKernelObjectParam() { return Params.front(); }
};

class SyclKernelBodyCreator
: public SyclKernelFieldHandler<SyclKernelBodyCreator> {
SyclKernelDeclCreator &DeclCreator;
llvm::SmallVector<Stmt *, 16> BodyStmts;
llvm::SmallVector<Stmt *, 16> FinalizeStmts;
llvm::SmallVector<Expr *, 16> InitExprs;
VarDecl *KernelObjClone;
InitializedEntity VarEntity;
CXXRecordDecl *KernelObj;
llvm::SmallVector<Expr *, 16> MemberExprBases;
FunctionDecl *KernelCallerFunc;
Expand All @@ -1029,22 +1061,24 @@ class SyclKernelBodyCreator
// statements in advance to allocate it, so we cannot do this as we go along.
CompoundStmt *createKernelBody() {

Expr *ILE = new (SemaRef.getASTContext()) InitListExpr(
SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation());
ILE->setType(QualType(KernelObj->getTypeForDecl(), 0));
KernelObjClone->setInit(ILE);
Stmt *FunctionBody = KernelCallerFunc->getBody();

ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin());
// Kernel object parameter from kernel caller function
ParmVarDecl *KernelCallerObjParam = *(KernelCallerFunc->param_begin());
// Kernel object parameter from generated kernel.
ParmVarDecl *KernelObjParam = DeclCreator.getKernelObjectParam();

// DeclRefExpr with valid source location but with decl which is not marked
// as used is invalid.
KernelObjClone->setIsUsed();
KernelObjParam->setIsUsed();
std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair =
std::make_pair(KernelObjParam, KernelObjClone);
std::make_pair(KernelCallerObjParam, KernelObjParam);

// Push the Kernel function scope to ensure the scope isn't empty
SemaRef.PushFunctionScope();

// Replacing all references to kernel caller function parameter in kernel
// body with references to kernel object parameter in generated kernel.
KernelBodyTransform KBT(MappingPair, SemaRef);
Stmt *NewBody = KBT.TransformStmt(FunctionBody).get();
BodyStmts.push_back(NewBody);
Expand Down Expand Up @@ -1094,26 +1128,13 @@ class SyclKernelBodyCreator
return Result;
}

// TODO: Correct Stream + Accessors
void createExprForStructOrScalar(FieldDecl *FD) {
ParmVarDecl *KernelParameter =
DeclCreator.getParamVarDeclsForCurrentField()[0];
InitializedEntity Entity =
InitializedEntity::InitializeMember(FD, &VarEntity);
QualType ParamType = KernelParameter->getOriginalType();
Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
SourceLocation());
if (FD->getType()->isPointerType() &&
FD->getType()->getPointeeType().getAddressSpace() !=
ParamType->getPointeeType().getAddressSpace())
DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(),
CK_AddressSpaceConversion, DRE, nullptr,
VK_RValue);
InitializationKind InitKind =
InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE);

ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE);
InitExprs.push_back(MemberInit.get());
}

void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base,
Expand Down Expand Up @@ -1153,32 +1174,8 @@ class SyclKernelBodyCreator
BodyStmts.push_back(Call);
}

// FIXME Avoid creation of kernel obj clone.
// See https://github.com/intel/llvm/issues/1544 for details.
static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC,
CXXRecordDecl *KernelObj) {
TypeSourceInfo *TSInfo =
KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr;
VarDecl *VD = VarDecl::Create(
Ctx, DC, SourceLocation(), SourceLocation(), KernelObj->getIdentifier(),
QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None);

return VD;
}

void handleSpecialType(FieldDecl *FD, QualType Ty) {
const auto *RecordDecl = Ty->getAsCXXRecordDecl();
// Perform initialization only if it is field of kernel object
if (MemberExprBases.size() == 1) {
InitializedEntity Entity =
InitializedEntity::InitializeMember(FD, &VarEntity);
// Initialize with the default constructor.
InitializationKind InitKind =
InitializationKind::CreateDefault(SourceLocation());
InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
InitExprs.push_back(MemberInit.get());
}
createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName,
FD);
}
Expand All @@ -1188,20 +1185,8 @@ class SyclKernelBodyCreator
CXXRecordDecl *KernelObj,
FunctionDecl *KernelCallerFunc)
: SyclKernelFieldHandler(S), DeclCreator(DC),
KernelObjClone(createKernelObjClone(S.getASTContext(),
DC.getKernelDecl(), KernelObj)),
VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)),
KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) {
markParallelWorkItemCalls();

Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
SourceLocation(), SourceLocation());
BodyStmts.push_back(DS);
DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create(
S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone,
false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0),
VK_LValue);
MemberExprBases.push_back(KernelObjCloneRef);
}

~SyclKernelBodyCreator() {
Expand Down Expand Up @@ -1237,15 +1222,40 @@ class SyclKernelBodyCreator
}

void handlePointerType(FieldDecl *FD, QualType FieldTy) final {
createExprForStructOrScalar(FD);
ParmVarDecl *KernelParameter =
DeclCreator.getParamVarDeclsForCurrentField()[0];
QualType ParamType = KernelParameter->getOriginalType();
Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
SourceLocation());
if (FD->getType()->isPointerType() &&
FD->getType()->getPointeeType().getAddressSpace() !=
ParamType->getPointeeType().getAddressSpace())
DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(),
CK_AddressSpaceConversion, DRE, nullptr,
VK_RValue);

MemberExpr *KernelObjectPointerField =
BuildMemberExpr(MemberExprBases.back(), FD);
Expr *AssignPointerParameter = BinaryOperator::Create(
SemaRef.getASTContext(), KernelObjectPointerField, DRE, BO_Assign,
FieldTy, VK_LValue, OK_Ordinary, SourceLocation(),
FPOptions(SemaRef.getASTContext().getLangOpts()));

BodyStmts.push_back(AssignPointerParameter);
}

// TODO: Accessors in structs
void handleStructType(FieldDecl *FD, QualType FieldTy) final {
createExprForStructOrScalar(FD);
// createExprForStructOrScalar(FD);
}

void handleScalarType(FieldDecl *FD, QualType FieldTy) final {
createExprForStructOrScalar(FD);
void handleKernelObject(CXXRecordDecl *KernelObject, QualType KernelType) {
ParmVarDecl *KernelParameter =
DeclCreator.getParamVarDeclsForCurrentField()[0];
QualType ParamType = KernelParameter->getOriginalType();
Expr *KernelObjRef = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType,
VK_LValue, SourceLocation());
MemberExprBases.push_back(KernelObjRef);
}

void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final {
Expand Down Expand Up @@ -1376,6 +1386,14 @@ class SyclKernelIntHeaderCreator
CurStruct = FD->getType()->getAsCXXRecordDecl();
CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8;
}
void handleKernelObject(CXXRecordDecl *KernelObject,
QualType KernelType) final {
uint64_t Size =
SemaRef.getASTContext().getTypeSizeInChars(KernelType).getQuantity();
// Offset for kernel object is 0
Header.addParamDesc(SYCLIntegrationHeader::kind_std_layout,
static_cast<unsigned>(Size), 0);
}

void leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD) final {
CurStruct = RD;
Expand Down Expand Up @@ -1447,7 +1465,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
StableName);

ConstructingOpenCLKernel = true;
VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body,
VisitRecordFields(KernelLambda, checker, kernel_decl, kernel_body,
int_header);
ConstructingOpenCLKernel = false;
}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/device-functions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ int main() {
kernel_single_task<class fake_kernel>([]() { foo(); });
return 0;
}
// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel()
// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject)
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %this)
// CHECK: define spir_func void @_Z3foov()
// CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/image_accessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,27 +7,27 @@
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
//
// CHECK-1DRO: %opencl.image1d_ro_t = type opaque
// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 8 %_arg_kernelObject, %opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-2DRO: %opencl.image2d_ro_t = type opaque
// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.0"* byval(%"class.{{.*}}.anon.0") align 8 %_arg_kernelObject, %opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-3DRO: %opencl.image3d_ro_t = type opaque
// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DRO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.1"* byval(%"class.{{.*}}.anon.1") align 8 %_arg_kernelObject, %opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-1DWO: %opencl.image1d_wo_t = type opaque
// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.2"* byval(%"class.{{.*}}.anon.2") align 8 %_arg_kernelObject, %opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-2DWO: %opencl.image2d_wo_t = type opaque
// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.3"* byval(%"class.{{.*}}.anon.3") align 8 %_arg_kernelObject, %opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-3DWO: %opencl.image3d_wo_t = type opaque
// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DWO: define spir_kernel void @{{.*}}(%"class.{{.*}}.anon.4"* byval(%"class.{{.*}}.anon.4") align 8 %_arg_kernelObject, %opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
//
// TODO: Add tests for the image_array opencl datatype support.
Expand Down
7 changes: 6 additions & 1 deletion clang/test/CodeGenSYCL/integration_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
// RUN: FileCheck -input-file=%t.h %s
//
// FIXME: Check incorrect header generation for accessor in base classes.
// XFAIL: *
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
//
// CHECK: class first_kernel;
Expand All @@ -27,22 +28,26 @@
// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-NEXT: //--- _ZTSZ4mainE12first_kernel
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 16 },
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 32 },
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTSN16second_namespace13second_kernelIcEE
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 },
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 },
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 16, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
// CHECK-EMPTY:
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,8 @@ void bar() {
[]() [[intelfpga::no_global_work_offset(0)]]{});
}

// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name1(%class.{{.*}}.Foo* byval(%class.{{.*}}.Foo) align 1 %_arg_kernelObject) {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name2(%"class.{{.*}}.anon"* byval(%"class.{{.*}}.anon") align 1 %_arg_kernelObject) {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK: define spir_kernel void @{{.*}}kernel_name3(%"class.{{.*}}.anon.0"* byval(%"class.{{.*}}.anon.0") align 1 %_arg_kernelObject) {{.*}} ![[NUM4:[0-9]+]]
// CHECK-NOT: ![[NUM4]] = !{i32 0}
// CHECK: ![[NUM5]] = !{}
Loading