Skip to content

[SYCL][Sema] Prevent Ctor/Dtor from being called during SYCL kernel r… #1

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

Open
wants to merge 1 commit into
base: sycl
Choose a base branch
from
Open

[SYCL][Sema] Prevent Ctor/Dtor from being called during SYCL kernel r… #1

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
9 changes: 8 additions & 1 deletion clang/lib/AST/Expr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -270,7 +270,14 @@ bool Expr::isFlexibleArrayMemberLike(
continue;
}
if (ConstantArrayTypeLoc CTL = TL.getAs<ConstantArrayTypeLoc>()) {
const Expr *SizeExpr = dyn_cast<IntegerLiteral>(CTL.getSizeExpr());
// FIXME: changed dyn_cast to dyn_cast_or_null
// to work around the fact that CTL.getSizeExpr() isn't set
// for a FieldDecl of a class generated from a lambda capture.
// This is highlighted only by the way lambda expression used
// as a SYCL kernel is being processed.
// In normal situation the capture list is used.
// No harm done, just a work around.
const Expr *SizeExpr = dyn_cast_or_null<IntegerLiteral>(CTL.getSizeExpr());
if (!SizeExpr || SizeExpr->getExprLoc().isMacroID())
return false;
}
Expand Down
462 changes: 233 additions & 229 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop

Large diffs are not rendered by default.

3 changes: 3 additions & 0 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
#pragma once

void* operator new (__SIZE_TYPE__ size, void* ptr) noexcept;
void* operator new[](__SIZE_TYPE__ size, void* ptr) noexcept;

#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
#define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]

Expand Down
2 changes: 2 additions & 0 deletions clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,8 @@
//
// Test which verifies that readonly attribute is generated for unexpected access mode value.

void* operator new (__SIZE_TYPE__ size, void* ptr) noexcept;

// Dummy library with unexpected access::mode enum value.
namespace sycl {
inline namespace _V1 {
Expand Down
40 changes: 25 additions & 15 deletions clang/test/CodeGenSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,11 @@ int main() {
// CHECK: [[ARG_C]].addr.ascast = addrspacecast ptr [[ARG_C]].addr to ptr addrspace(4)
//
// Lambda object alloca
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
// CHECK: [[KERNEL_UNION_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[KERNEL]] to %union.__wrapper_union addrspace(4)*
// CHECK: [[KERNEL_OBJ_PTR_ALLOCA:%[a-zA-Z0-9_.]+]] = alloca %class{{.*}}.anon addrspace(4)*, align 8
// CHECK: [[KERNEL_OBJ_PTR:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon addrspace(4)** [[KERNEL_OBJ_PTR_ALLOCA]] to %class.anon addrspace(4)* addrspace(4)*

//
// Kernel argument stores
// CHECK: store i32 [[ARG_A]], ptr addrspace(4) [[ARG_A]].addr.ascast
Expand All @@ -60,13 +64,15 @@ int main() {
// CHECK: store i32 [[ARG_C]], ptr addrspace(4) [[ARG_C]].addr.ascast
//
// Check A and B scalar fields initialization
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 0
// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_A]].addr.ascast
// CHECK: store i32 [[ARG_A_LOAD]], ptr addrspace(4) [[FIELD_A]]
// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 1
// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_B]].addr.ascast
// CHECK: store i32 [[ARG_B_LOAD]], ptr addrspace(4) [[FIELD_B]]
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_UNION_OBJ]] to %class{{.*}}.anon addrspace(4)*
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
// CHECK: store i32 [[ARG_A_LOAD]], i32 addrspace(4)* [[FIELD_A]]
// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1
// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_B]].addr.ascast
// CHECK: store i32 [[ARG_B_LOAD]], i32 addrspace(4)* [[FIELD_B]]
//
// Check accessors initialization
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 2
Expand All @@ -82,11 +88,15 @@ int main() {
// CHECK: store i32 [[ARG_C_LOAD]], ptr addrspace(4) [[FIELD_C]]
//
// Check __init method calls
// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP2]], i32 0, i32 2
// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC1_DATA]].addr.ascast
// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[ACC1_FIELD]], ptr addrspace(1) noundef [[ACC1_DATA_LOAD]]
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[KERNEL_OBJ_PTR]]
// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
//
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC2_DATA]].addr.ascast
// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}}, ptr addrspace(1) noundef [[ACC2_DATA_LOAD]]
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[KERNEL_OBJ_PTR]]
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
// CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]
4 changes: 4 additions & 0 deletions clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,8 @@
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s

// Set the cursor so that CEHCK-DAG don't pickup wrong matches
// CHECK: define {{.*}} @_Z6usagesv

void bar(int & Data) {}
// CHECK-DAG: define {{.*}}spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
void bar2(int & Data) {}
Expand Down
8 changes: 6 additions & 2 deletions clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,10 @@ int main() {
// Check alloca for pointer argument
// CHECK: [[MEM_ARG]].addr = alloca ptr addrspace(1)
// Check lambda object alloca
// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
// CHECK: [[UNIONALLOCA:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
// CHECK: [[UNION:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[UNIONALLOCA]] to %union.__wrapper_union addrspace(4)*
// CHECK: [[ANONPTRALLOCA:%[a-zA-Z0-9_.]+]] = alloca %class{{.*}}.anon addrspace(4)*, align 8
// CHECK: [[ANONPTRALLOCA_PTR:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon addrspace(4)** [[ANONPTRALLOCA]] to %class.anon addrspace(4)* addrspace(4)*
// Check allocas for ranges
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
Expand All @@ -44,7 +47,8 @@ int main() {
// CHECK: call spir_func {{.*}}accessor

// Check accessor GEP
// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[ANON]], i32 0, i32 0
// CHECK: [[ANON:%[0-9]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[ANONPTRALLOCA_PTR]]
// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON]], i32 0, i32 0

// Check load from kernel pointer argument alloca
// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG]].addr.ascast
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -317,20 +317,20 @@ int main() {

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
// CHECK-NOT: noalias
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo8v()
Functor10 f10;
h.single_task<class kernel_name32>(f10);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
// CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
Foo8 boo8;
h.single_task<class kernel_name33>(boo8);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #3 align 2
h.single_task<class kernel_name34>(
[]() [[intel::kernel_args_restrict]]{});

Expand Down
18 changes: 13 additions & 5 deletions clang/test/CodeGenSYCL/generated-types-initialization.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,8 +41,10 @@ int main() {
// CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
//
// Kernel object clone.
// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %union.__wrapper_union
// CHECK: %[[K_PTR_ALLOCA:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4)
// CHECK: %[[K_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[K]] to ptr addrspace(4)
// CHECK: %[[K_PTR_ALLOCA_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[K_PTR_ALLOCA]] to ptr addrspace(4)
//
// Argument reference.
// CHECK: %[[Arg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_Obj to ptr addrspace(4)
Expand All @@ -52,20 +54,26 @@ int main() {
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[GEP]], ptr addrspace(4) align 8 %[[Arg_ref]], i64 16, i1 false)
//
// Kernel body call.
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
// CHECK: store ptr addrspace(4) %[[K_as_cast]], ptr addrspace(4) %[[K_PTR_ALLOCA_as_cast]]
// CHECK: %[[K_PTR:[a-zA-Z0-9_.]+]] = load ptr addrspace(4), ptr addrspace(4) %[[K_PTR_ALLOCA_as_cast]]
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_PTR]])

// CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
//
// Kernel object clone.
// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2
// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %union.__wrapper_union.2
// CHECK: %[[NNSK_PTR_ALLOCA:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4)
// CHECK: %[[NNSK_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[NNSK]] to ptr addrspace(4)
// CHECK: %[[NNSK_PTR_ALLOCA_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[NNSK_PTR_ALLOCA]] to ptr addrspace(4)
//
// Argument reference.
// CHECK: %[[NNSArg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_NNSObj to ptr addrspace(4)
//
// Initialization.
// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.2, ptr addrspace(4) %[[NNSK_as_cast]], i32 0, i32 0
// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.3, ptr addrspace(4) %[[NNSK_as_cast]], i32 0, i32 0
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[NNSGEP]], ptr addrspace(4) align 8 %[[NNSArg_ref]], i64 16, i1 false)
//
// Kernel body call.
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[NNSK_as_cast]])
// CHECK: store ptr addrspace(4) %[[NNSK_as_cast]], ptr addrspace(4) %[[NNSK_PTR_ALLOCA_as_cast]]
// CHECK: %[[NNSK_PTR:[a-zA-Z0-9_.]+]] = load ptr addrspace(4), ptr addrspace(4) %[[NNSK_PTR_ALLOCA_as_cast]]
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[NNSK_PTR]])
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,22 +7,22 @@
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
//
// CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
//
// TODO: Add tests for the image_array opencl datatype support.
#include "Inputs/sycl.hpp"
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,24 +56,24 @@ int main() {

// Check allocas for kernel parameters and local functor object
// CHECK: %[[ARG_A_ALLOCA:[a-zA-Z0-9_.]+]] = alloca i32, align 4
// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8
// CHECK: %[[UNIONALLOCA:[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4), align 8
// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[ARG_A_ALLOCA]] to ptr addrspace(4)
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECT_ALLOCA]] to ptr addrspace(4)
// CHECK: %[[UNION:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[UNIONALLOCA]] to ptr addrspace(4)
// CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base to ptr addrspace(4)
// CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base1 to ptr addrspace(4)
// CHECK: store i32 %_arg_a, ptr addrspace(4) %[[ARG_A]], align 4

// Initialize 'base' subobject
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[LOCAL_OBJECT]], ptr addrspace(4) align 4 %[[ARG_BASE]], i64 12, i1 false)
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[UNION]], ptr addrspace(4) align 4 %[[ARG_BASE]], i64 12, i1 false)

// Initialize 'second_base' subobject
// First, derived-to-base cast with offset:
// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16
// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[UNION]], i64 16
// Initialize 'second_base'
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 24, i1 false)

// Initialize field 'a'
// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3
// CHECK: %[[LOAD_A:[0-9]+]] = load i32, ptr addrspace(4) %[[ARG_A]], align 4
// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[UNION]], i32 0, i32 3
// CHECK: store i32 %[[LOAD_A]], ptr addrspace(4) %[[GEP_A]]

Loading