Fix alignment of generated type (#8118)

Alignment can be explicitly specified on records/fields using attribute.
Copy attributes from original type to compiler generated types to
maintain same alignment in generated types

Author: elizabethandrews

clang/lib/Sema/SemaSYCL.cpp

@@ -1973,6 +1973,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);
   }
 
@@ -1987,6 +1989,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);
   }

clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp

@@ -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;
@@ -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)

clang/test/CodeGenSYCL/pointers-in-structs.cpp

@@ -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;
@@ -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)