Skip to content

Fix alignment of generated type #8118

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

Merged
merged 1 commit into from
Jan 31, 2023
Merged

Fix alignment of generated type #8118

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
4 changes: 4 additions & 0 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1994,6 +1994,8 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler {
const_cast<DeclContext *>(RD->getDeclContext()), SourceLocation(),
SourceLocation(), getModifiedName(RD->getIdentifier()));
ModifiedRD->startDefinition();
if (RD->hasAttrs())
ModifiedRD->setAttrs(RD->getAttrs());
ModifiedRecords.push_back(ModifiedRD);
}

Expand All @@ -2008,6 +2010,8 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler {
Ctx.getTrivialTypeSourceInfo(FieldTy, SourceLocation()), /*BW=*/nullptr,
/*Mutable=*/false, ICIS_NoInit);
Field->setAccess(FD->getAccess());
if (FD->hasAttrs())
Field->setAttrs(FD->getAttrs());
// Add generated field to generated record.
ModifiedRecords.back()->addDecl(Field);
}
Expand Down
27 changes: 21 additions & 6 deletions clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,26 +1,35 @@
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s

// This test checks that compiler generates correct address spaces for pointer
// kernel arguments that are wrapped by struct.
// kernel arguments that are wrapped by struct. Generated class should retain
// original padding and alignment.

#include "Inputs/sycl.hpp"

struct A {
float *F;
};

struct B {
struct alignas(16) B {
int *F1;
float *F2;
A F3;
int *F4[2];
};

struct testFieldAlignment {
int *ptr;
alignas(16) float arr[4];
int data;
};

int main() {
B Obj;
B Obj1;
testFieldAlignment Obj2;
sycl::kernel_single_task<class structs>(
[=]() {
(void)Obj;
(void)Obj1;
(void)Obj2;
});
float A = 1;
float *Ptr = &A;
Expand All @@ -33,9 +42,15 @@ int main() {
return 0;
}

// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)*, float addrspace(1)*, %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x i32 addrspace(1)*] }
// Padding in generated class class should match the 'original' padding
// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)*, float addrspace(1)*, %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x i32 addrspace(1)*], [8 x i8] }
// CHECK: %[[GENERATED_A]] = type { float addrspace(1)* }
// CHECK: %[[GENERATED_TESTFIELDALIGNMENT:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)*, [8 x i8], [4 x float], i32, [12 x i8] }
// CHECK: %struct.B = type { i32 addrspace(4)*, float addrspace(4)*, %struct.A, [2 x i32 addrspace(4)*], [8 x i8] }
// CHECK: %struct.A = type { float addrspace(4)* }
// %struct.testFieldAlignment = type { i32 addrspace(4)*, [8 x i8], [4 x float], i32, [12 x i8] }
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
// CHECK-SAME: %[[GENERATED_B]]* noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
// CHECK-SAME: %[[GENERATED_B]]* noundef byval(%[[GENERATED_B]]) align 16 %_arg_Obj1
// CHECK-SAME: %[[GENERATED_TESTFIELDALIGNMENT]]* noundef byval(%[[GENERATED_TESTFIELDALIGNMENT]]) align 16 %_arg_Obj2
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)
29 changes: 22 additions & 7 deletions clang/test/CodeGenSYCL/pointers-in-structs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,26 +1,35 @@
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s

// This test checks that compiler generates correct address spaces for pointer
// kernel arguments that are wrapped by struct.
// kernel arguments that are wrapped by struct. Generated class should retain
// original padding and alignment.

#include "Inputs/sycl.hpp"

struct A {
float *F;
};

struct B {
struct alignas(16) B {
int *F1;
float *F2;
A F3;
int *F4[2];
};

struct testFieldAlignment {
int *ptr;
alignas(16) float arr[4];
int data;
};

int main() {
B Obj;
B Obj1;
testFieldAlignment Obj2;
sycl::kernel_single_task<class structs>(
[=]() {
(void)Obj;
(void)Obj1;
(void)Obj2;
});
float A = 1;
float *Ptr = &A;
Expand All @@ -33,10 +42,16 @@ int main() {
return 0;
}

// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1), ptr addrspace(1), %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x ptr addrspace(1)] }
// CHECK: [[GENERATED_A]] = type { ptr addrspace(1) }
// Padding in generated class class should match the 'original' padding
// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1), ptr addrspace(1), %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x ptr addrspace(1)], [8 x i8] }
// CHECK: %[[GENERATED_A]] = type { ptr addrspace(1) }
// CHECK: %[[GENERATED_TESTFIELDALIGNMENT:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1), [8 x i8], [4 x float], i32, [12 x i8] }
// CHECK: %struct.B = type { ptr addrspace(4), ptr addrspace(4), %struct.A, [2 x ptr addrspace(4)], [8 x i8] }
// CHECK: %struct.testFieldAlignment = type { ptr addrspace(4), [8 x i8], [4 x float], i32, [12 x i8] }
// %struct.A = type { ptr addrspace(4) }
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }

// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
// CHECK-SAME: ptr noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
// CHECK-SAME: ptr noundef byval(%[[GENERATED_B]]) align 16 %_arg_Obj1
// CHECK-SAME: ptr noundef byval(%[[GENERATED_TESTFIELDALIGNMENT]]) align 16 %_arg_Obj2
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(ptr noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)