[SYCL] Add support for union types as kernel parameter #2285
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -830,6 +830,13 @@ class KernelObjVisitor { | |
| else if (ElementTy->isStructureOrClassType()) | ||
| VisitRecord(Owner, ArrayField, ElementTy->getAsCXXRecordDecl(), | ||
| handlers...); | ||
| else if (ElementTy->isUnionType()) | ||
| // TODO: This check is still necessary I think?! Array seems to handle | ||
| // this differently (see above) for structs I think. | ||
| // if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { | ||
| VisitUnion(Owner, ArrayField, ElementTy->getAsCXXRecordDecl(), | ||
| handlers...); | ||
| //} | ||
| else if (ElementTy->isArrayType()) | ||
| VisitArrayElements(ArrayField, ElementTy, handlers...); | ||
| else if (ElementTy->isScalarType()) | ||
|
|
@@ -857,6 +864,40 @@ class KernelObjVisitor { | |
| void VisitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, | ||
| const CXXRecordDecl *Wrapper, Handlers &... handlers); | ||
|
|
||
| // Base case, only calls these when filtered. | ||
| template <typename... FilteredHandlers, typename ParentTy> | ||
| void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, | ||
| const CXXRecordDecl *Wrapper, | ||
| FilteredHandlers &... handlers) { | ||
| (void)std::initializer_list<int>{ | ||
erichkeane marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| (handlers.enterUnion(Owner, Parent), 0)...}; | ||
| VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...); | ||
erichkeane marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| (void)std::initializer_list<int>{ | ||
| (handlers.leaveUnion(Owner, Parent), 0)...}; | ||
| } | ||
|
|
||
| template <typename... FilteredHandlers, typename ParentTy, | ||
| typename CurHandler, typename... Handlers> | ||
| std::enable_if_t<!CurHandler::VisitUnionBody> | ||
| VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, | ||
| const CXXRecordDecl *Wrapper, | ||
| FilteredHandlers &... filtered_handlers, CurHandler &cur_handler, | ||
| Handlers &... handlers) { | ||
| VisitUnion<FilteredHandlers...>(Owner, Parent, Wrapper, | ||
| filtered_handlers..., handlers...); | ||
| } | ||
|
|
||
| template <typename... FilteredHandlers, typename ParentTy, | ||
| typename CurHandler, typename... Handlers> | ||
| std::enable_if_t<CurHandler::VisitUnionBody> | ||
| VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, | ||
| const CXXRecordDecl *Wrapper, | ||
| FilteredHandlers &... filtered_handlers, CurHandler &cur_handler, | ||
| Handlers &... handlers) { | ||
| VisitUnion<FilteredHandlers..., CurHandler>( | ||
| Owner, Parent, Wrapper, filtered_handlers..., cur_handler, handlers...); | ||
| } | ||
|
|
||
| template <typename... Handlers> | ||
| void VisitRecordHelper(const CXXRecordDecl *Owner, | ||
| clang::CXXRecordDecl::base_class_const_range Range, | ||
|
|
@@ -942,6 +983,11 @@ class KernelObjVisitor { | |
| CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); | ||
| VisitRecord(Owner, Field, RD, handlers...); | ||
| } | ||
| } else if (FieldTy->isUnionType()) { | ||
| if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { | ||
| CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); | ||
| VisitUnion(Owner, Field, RD, handlers...); | ||
| } | ||
| } else if (FieldTy->isReferenceType()) | ||
| KF_FOR_EACH(handleReferenceType, Field, FieldTy); | ||
| else if (FieldTy->isPointerType()) | ||
|
|
@@ -1005,6 +1051,7 @@ class SyclKernelFieldHandler { | |
| } | ||
| virtual bool handleSyclHalfType(FieldDecl *, QualType) { return true; } | ||
| virtual bool handleStructType(FieldDecl *, QualType) { return true; } | ||
| virtual bool handleUnionType(FieldDecl *, QualType) { return true; } | ||
| virtual bool handleReferenceType(FieldDecl *, QualType) { return true; } | ||
| virtual bool handlePointerType(FieldDecl *, QualType) { return true; } | ||
| virtual bool handleArrayType(FieldDecl *, QualType) { return true; } | ||
|
|
@@ -1024,6 +1071,8 @@ class SyclKernelFieldHandler { | |
| virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) { | ||
| return true; | ||
| } | ||
| virtual bool enterUnion(const CXXRecordDecl *, FieldDecl *) { return true; } | ||
| virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; } | ||
|
|
||
| // The following are used for stepping through array elements. | ||
|
|
||
|
|
@@ -1201,11 +1250,71 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { | |
| } | ||
| }; | ||
|
|
||
| // A type to check the validity of passing union with accessor/sampler/stream | ||
| // member as a kernel argument types. | ||
| class SyclKernelUnionBodyChecker : public SyclKernelFieldHandler { | ||
| static constexpr const bool VisitUnionBody = true; | ||
|
There was a problem hiding this comment. I think your problem is that SyclKernelFieldHandler needs this line as well, except =false instead. I think that results in the union-filtering not actually doing anything. There was a problem hiding this comment.
There was a problem hiding this comment. I can't find where it is done, that might be the reason why everything is decomposed. There was a problem hiding this comment. Yep, it should be in the base class, not just the other checker. WHich is why the tests are still wrong. |
||
| int UnionCount = 0; | ||
| bool IsInvalid = false; | ||
| DiagnosticsEngine &Diag; | ||
|
|
||
| public: | ||
| SyclKernelUnionBodyChecker(Sema &S) | ||
| : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} | ||
| bool isValid() { return !IsInvalid; } | ||
|
|
||
| bool enterUnion(const CXXRecordDecl *RD, FieldDecl *FD) { | ||
| ++UnionCount; | ||
| return true; | ||
| } | ||
|
|
||
| bool leaveUnion(const CXXRecordDecl *RD, FieldDecl *FD) { | ||
| --UnionCount; | ||
| return true; | ||
| } | ||
|
|
||
| bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { | ||
erichkeane marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| if (UnionCount) { | ||
|
There was a problem hiding this comment. This is a NIT, so don't redo the patch just for this (but please do if you have to push another patch), but I'd prefer the body of this to be extracted into a private function. Something like: Then, all of these other functions simply become:
I want this, since I believe it will make the 'printing the diagnostic trail' much easier. |
||
| IsInvalid = true; | ||
| Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) | ||
|
There was a problem hiding this comment. You need a new diagnostic that says about union members. This one is incorrect here. |
||
| << FieldTy; | ||
| } | ||
| return isValid(); | ||
| } | ||
|
|
||
| bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { | ||
| if (UnionCount) { | ||
| IsInvalid = true; | ||
| Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) | ||
| << FieldTy; | ||
| } | ||
| return isValid(); | ||
| } | ||
|
|
||
| bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { | ||
bader marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| if (UnionCount) { | ||
| IsInvalid = true; | ||
| Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) | ||
| << FieldTy; | ||
| } | ||
| return isValid(); | ||
| } | ||
| bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { | ||
| if (UnionCount) { | ||
| IsInvalid = true; | ||
| Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) | ||
| << FieldTy; | ||
| } | ||
| return isValid(); | ||
| } | ||
| }; | ||
|
|
||
| // A type to Create and own the FunctionDecl for the kernel. | ||
| class SyclKernelDeclCreator : public SyclKernelFieldHandler { | ||
| FunctionDecl *KernelDecl; | ||
| llvm::SmallVector<ParmVarDecl *, 8> Params; | ||
| SyclKernelFieldChecker &ArgChecker; | ||
| SyclKernelUnionBodyChecker &ArgChecker1; | ||
|
There was a problem hiding this comment. I'd probably call this 'UnionChecker'. Additionally, note that #2289 makes this part unnecessary. There was a problem hiding this comment. Ok. I will change that. There was a problem hiding this comment. Actually, looking again, adding ArgChecker1 is completely unnecessary. ArgChecker itself is completely unnecessary here too (and likely should be removed). I remove it in #2289 (depending on when that gets in), but it seems to be a leftover from when it was necessary. There was a problem hiding this comment. It seems like only place we can add this handler "SyclKernelUnionBodyChecker checker1(*this)" to ConstructOpenCLKernel. There was a problem hiding this comment. Yep, I thought you would have noticed trying to debug why diagnostics weren't working right, but it needs to be added to these two lines instead: Visitor.VisitRecordBases(KernelObj, checker, kernel_decl, kernel_body, Visitor.VisitRecordFields(KernelObj, checker, kernel_decl, kernel_body, |
||
| Sema::ContextRAII FuncContext; | ||
| // Holds the last handled field's first parameter. This doesn't store an | ||
| // iterator as push_back invalidates iterators. | ||
|
|
@@ -1340,12 +1449,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { | |
|
|
||
| public: | ||
| SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, | ||
| SyclKernelUnionBodyChecker &ArgChecker1, StringRef Name, | ||
| SourceLocation Loc, bool IsInline, bool IsSIMDKernel) | ||
| : SyclKernelFieldHandler(S), | ||
| KernelDecl(createKernelDecl(S.getASTContext(), Name, Loc, IsInline, | ||
| IsSIMDKernel)), | ||
| ArgChecker(ArgChecker), ArgChecker1(ArgChecker1), | ||
| FuncContext(SemaRef, KernelDecl) {} | ||
|
|
||
| ~SyclKernelDeclCreator() { | ||
| ASTContext &Ctx = SemaRef.getASTContext(); | ||
|
|
@@ -1416,6 +1526,10 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { | |
| return true; | ||
| } | ||
|
|
||
| bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { | ||
| return handleScalarType(FD, FieldTy); | ||
| } | ||
|
|
||
| bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final { | ||
| addParam(FD, FieldTy); | ||
| return true; | ||
|
|
@@ -1751,6 +1865,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { | |
| return true; | ||
| } | ||
|
|
||
| bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { | ||
| return handleScalarType(FD, FieldTy); | ||
| } | ||
|
|
||
| bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { | ||
| CXXCastPath BasePath; | ||
| QualType DerivedTy(RD->getTypeForDecl(), 0); | ||
|
|
@@ -1955,6 +2073,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { | |
| return true; | ||
| } | ||
|
|
||
| bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { | ||
| return handleScalarType(FD, FieldTy); | ||
| } | ||
|
|
||
| bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { | ||
| addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); | ||
| return true; | ||
|
|
@@ -2057,8 +2179,9 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, | |
| Diag(LC.getLocation(), diag::err_implicit_this_capture); | ||
| } | ||
| SyclKernelFieldChecker checker(*this); | ||
| SyclKernelUnionBodyChecker checker1(*this); | ||
| SyclKernelDeclCreator kernel_decl( | ||
| *this, checker, checker1, KernelName, KernelObj->getLocation(), | ||
| KernelCallerFunc->isInlined(), KernelCallerFunc->hasAttr<SYCLSimdAttr>()); | ||
| SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj, | ||
| KernelCallerFunc); | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,70 @@ | ||
| // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out | ||
| // RUN: FileCheck -input-file=%t.h %s | ||
|
|
||
| // This test checks the integration header generated when | ||
| // the kernel argument is union. | ||
|
|
||
| // CHECK: #include <CL/sycl/detail/kernel_desc.hpp> | ||
|
|
||
| // CHECK: class MyKernel; | ||
|
|
||
| // CHECK: __SYCL_INLINE_NAMESPACE(cl) { | ||
| // CHECK-NEXT: namespace sycl { | ||
| // CHECK-NEXT: namespace detail { | ||
|
|
||
| // CHECK: static constexpr | ||
| // CHECK-NEXT: const char* const kernel_names[] = { | ||
| // CHECK-NEXT: "_ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel" | ||
| // CHECK-NEXT: }; | ||
|
|
||
| // CHECK: static constexpr | ||
| // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { | ||
| // CHECK-NEXT: //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 12, 16 }, | ||
|
There was a problem hiding this comment. I don't think this is right, is it? 'S' is just the union, right? We should be capturing that as a single std_layout, not as however many these are. |
||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, | ||
| // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, | ||
| // CHECK-EMPTY: | ||
| // CHECK-NEXT:}; | ||
|
|
||
| // CHECK: static constexpr | ||
| // CHECK-NEXT: const unsigned kernel_signature_start[] = { | ||
| // CHECK-NEXT: 0 // _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel | ||
| // CHECK-NEXT: }; | ||
|
|
||
| // CHECK: template <> struct KernelInfo<class MyKernel> { | ||
|
|
||
| #include "sycl.hpp" | ||
|
|
||
| using namespace cl::sycl; | ||
|
|
||
| union MyNestedUnion { | ||
| int FldArr[1]; | ||
| float FldFloat; | ||
| }; | ||
|
|
||
| union MyUnion { | ||
| int FldInt; | ||
| MyNestedUnion FldUnion; | ||
| int FldArr[3]; | ||
| }; | ||
|
|
||
| MyUnion GlobS; | ||
|
|
||
| bool test0() { | ||
| MyUnion S = GlobS; | ||
| MyUnion S0 = {0}; | ||
| { | ||
| buffer<MyUnion, 1> Buf(&S0, range<1>(1)); | ||
| queue myQueue; | ||
| myQueue.submit([&](handler &cgh) { | ||
| auto B = Buf.get_access<access::mode::write>(cgh); | ||
| cgh.single_task<class MyKernel>([=] { B; S; }); | ||
| }); | ||
| } | ||
| } | ||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,83 @@ | ||
| // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // This test checks if compiler accepts union as kernel parameters. | ||
|
|
||
| #include "sycl.hpp" | ||
|
|
||
| using namespace cl::sycl; | ||
|
|
||
| union MyUnion { | ||
| int FldInt; | ||
| int FldArr[3]; | ||
| }; | ||
|
|
||
| MyUnion GlobS; | ||
|
|
||
| void test0() { | ||
| MyUnion S = GlobS; | ||
| MyUnion S0 = {0}; | ||
| { | ||
| buffer<MyUnion, 1> Buf(&S0, range<1>(1)); | ||
| queue myQueue; | ||
| myQueue.submit([&](handler &cgh) { | ||
| auto B = Buf.get_access<access::mode::write>(cgh); | ||
| cgh.single_task<class MyKernel>([=] {B; S; }); | ||
| }); | ||
| } | ||
| } | ||
|
|
||
| // CHECK MyKernel parameters | ||
|
There was a problem hiding this comment. @elizabethandrews can you take a look at this? The union should be a simple un-decomposed value. There was a problem hiding this comment. From the Sema tests below, it looks like parameters are being passed individually as well? This in incorrect right? There was a problem hiding this comment. Yes, the test seems incorrect. I will update the test. Thanks Erich and Elizabeth. |
||
| // CHECK: define spir_kernel void @{{.*}}MyKernel | ||
| // CHECK-SAME: %union.{{.*}}.MyUnion addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+1]], | ||
| // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_ARG3:%[a-zA-Z0-9_]+2]], | ||
| // CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]], | ||
| // CHECK-SAME: %union.{{.*}}.MyUnion* byval(%union.{{.*}}.MyUnion) align 4 [[MEM_ARG4:%[a-zA-Z0-9_]+4]], | ||
| // CHECK-SAME: i32 [[MEM1:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: i32 [[MEM2_Array1:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: i32 [[MEM2_Array2:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: i32 [[MEM2_Array3:%[a-zA-Z0-9_]+]] | ||
|
|
||
| // Check alloca for pointer arguments | ||
| // CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca %union._ZTS7MyUnion.MyUnion addrspace(1)*, align 8 | ||
| // CHECK: [[MEM1]].addr{{[0-9]*}} = alloca i32, align 4 | ||
| // CHECK: [[MEM2_Array1]].addr{{[0-9]*}} = alloca i32, align 4 | ||
| // CHECK: %_arg_FldArr.addr6 = alloca i32, align 4 | ||
| // CHECK: %_arg_FldArr.addr8 = alloca i32, align 4 | ||
|
|
||
| // Check lambda object alloca | ||
| // CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4 | ||
|
|
||
| // Check allocas for ranges | ||
| // CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" | ||
| // CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" | ||
| // CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" | ||
|
|
||
| // CHECK: store i32 [[MEM1]], i32* [[MEM1]].addr{{[0-9]*}}, align 4 | ||
| // CHECK: store i32 [[MEM2_Array1]], i32* [[MEM2_Array1]].addr{{[0-9]*}}, align 4 | ||
| // CHECK: store i32 [[MEM2_Array2]], i32* %_arg_FldArr.addr6, align 4 | ||
| // CHECK: store i32 [[MEM2_Array3]], i32* %_arg_FldArr.addr8, align 4 | ||
| // | ||
| // CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 | ||
| // CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[L_STRUCT_ADDR]] to %"class{{.*}}accessor" addrspace(4)* | ||
| // CHECK: call spir_func void @{{.*}}MyUnion{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]]) | ||
| // CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 1 | ||
| // CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion* [[Z0]] to i8* | ||
| // CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion* [[MEM_ARG4]] to i8* | ||
| // CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 12, i1 false) | ||
| // CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 | ||
|
|
||
| // Check load from kernel pointer argument alloca | ||
| // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load %union._ZTS7MyUnion.MyUnion addrspace(1)*, %union._ZTS7MyUnion.MyUnion addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}, align 8 | ||
| // CHECK: [[MEMCPY_DST1:%[0-9a-zA-Z_]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE1]] to i8* | ||
| // CHECK: [[MEMCPY_SRC1:%[0-9a-zA-Z_]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_ARG2]] to i8* | ||
| // call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST1]], i8* align 4 [[MEMCPY_SRC1]], i64 4, i1 false), !tbaa.struct [[ACC_CAST2:%[0-9]+]] | ||
| // CHECK: [[MEMCPY_DST2:%[0-9a-zA-Z_]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE2]] to i8* | ||
| // CHECK: [[MEMCPY_SRC2:%[0-9a-zA-Z_]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_ARG3]] to i8* | ||
| // call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST2]], i8* align 4 [[MEMCPY_SRC2]], i64 4, i1 false), !tbaa.struct [[ACC_CAST2:%[0-9]+]] | ||
|
|
||
| // Check __init method call | ||
| // CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)* | ||
| // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], %union._ZTS7MyUnion.MyUnion addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) | ||
| // CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}.anon"* [[LOCAL_OBJECT]] to %"class{{.*}}.anon" addrspace(4)* | ||
| // CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* [[ACC_CAST2]]) | ||
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.