[OpenCL] Mark kernel arguments as ABI aligned

[nikic]

Summary

Following the discussion on D118229, this marks all pointer-typed kernel arguments as having ABI alignment, hopefully addressing the regression from that change.

[yaxunl]

LGTM. I will defer to Anastasia.

clang/lib/CodeGen/CGCall.cpp

2488

Pls add a reference to the OpenCL spec in the comment:

OpenCL spec v3.0.10 Appendix C Alignment of Application Data Types:
The user is responsible for ensuring that pointers passed into and out of OpenCL kernels are natively aligned relative to the data type of the parameter as defined in the kernel language and SPIR-V specifications.

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2489

Why does opencl need special treatment here? Isn't this the default behavior for all abi/languages?

[nikic]

Add spec quote.

[nikic]

clang/lib/CodeGen/CGCall.cpp

2489

In C++, references are required to be aligned (see the case directly above this), but pointers generally aren't.

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2489

Plus I would assume this also applies to non-kernels

2489

But pre-opaque pointers, these would have ended up using the datalayout ABI alignment?

[nikic]

clang/lib/CodeGen/CGCall.cpp

2489

No, at least not in generic code. We used to assume this, but it has been fixed a few releases back.

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

2489

I agree it doesn't seem to be kernel-specific, should this apply to all functions? However, I wonder why this should need a special case in OpenCL C compared to C behavior?

I feel the spec reference is not entirely related as all it's trying to say is that the host code is responsible to make sure the data are aligned in accordance with the device.

Looking at LangRef it doesn't state what the default values for alignments are:
https://llvm.org/docs/LangRef.html#parameter-attributes
but the change does seem to make sense in principle.

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

2489

Ok I think this might be a more relevant reference:
https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#alignment-of-types

For arguments to a __kernel function declared to be a pointer to a data type, the OpenCL compiler can assume that the pointee is always appropriately aligned as required by the data type.

However it still feels like it says that we expect data passed from host are adequately aligned?

So I feel like this might be an implementation specific aspect that spec doesn't actually regulate...

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2489

The conformance test does check the low bits of the pointer are aligned

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2489

I don't think in general we can assume pointers are naturally aligned. There are valid use cases of pointers to non-aligned types, e.g. a member of a packed structure. A function with a pointer-type argument should accept pointer arguments to non-aligned types. OpenCL pointer-type kernel arguments are special cases where the alignment is guaranteed by the spec.

@rjmccall

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

2489

There are valid use cases of pointers to non-aligned types, e.g. a member of a packed structure.

What if such types are used in kernel arguments, I can't see any related restrictions...

The conformance test does check the low bits of the pointer are aligned

Ok, just to understand a bit more... has the failure been introduced recently or is this a new test?

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2489

The test started failing recently due to no longer using the pointee element types' ABI alignment. clang needs to start providing it in the IR explicitly

[nikic]

clang/lib/CodeGen/CGCall.cpp

2489

There are valid use cases of pointers to non-aligned types, e.g. a member of a packed structure.

What if such types are used in kernel arguments, I can't see any related restrictions...

Isn't this covered by the part of the OpenCL spec you quoted above? Pointers can generally be unaligned, but for pointer arguments to kernels the spec requires that they are aligned.

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

2489

I am still confused why the fix is only needed for the kernel functions now... Has there been any change in LLVM IR that triggers this issue? Is there some specific handling happening in for the functions that are regular and therefore are being called in contrast to kernels that are normally not called on the kernel side?

[Anastasia]

clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl

3

Let's choose the pointee type with width different from pointer size... also should we test types that @yaxunl mentioned e.g with packed structs or void?

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

2489

To unconfuse (mainly myself) the difference here is that OpenCL spec actually fixes the alignment of types in contrast to C/C++ that only specifies the alignment requirements. However OpenCL spec does special case kernel function parameters which is why the normal functions still remain regulated by C/C++ rules... not sure it was intentional but the change is reasonable.

[nikic]

Update spec reference, add a dedicated alignment test.

[arsenm]

clang/test/CodeGenOpenCL/kernel-param-alignment.cl

11–14

Should throw in some vectors for good measure

[nikic]

Check vector alignment.

clang/lib/CodeGen/CGCall.cpp

if (const auto *RefTy = ParamType->getAs<ReferenceType>()) {
         Attrs.addAttribute(llvm::Attribute::NonNull);
       if (PTy->isObjectType()) {
         llvm::Align Alignment =
             getNaturalPointeeTypeAlignment(ParamType).getAsAlign();
         Attrs.addAlignmentAttr(Alignment);
       }
     }
 
+    // From OpenCL spec v3.0.10 section 6.3.5 Alignment of Types:
+    // > For arguments to a __kernel function declared to be a pointer to a
+    // > data type, the OpenCL compiler can assume that the pointee is always
+    // > appropriately aligned as required by the data type.
+    if (TargetDecl && TargetDecl->hasAttr<OpenCLKernelAttr>() &&
+        ParamType->isPointerType()) {
+      QualType PTy = ParamType->getPointeeType();
+      if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {
+        llvm::Align Alignment =
+            getNaturalPointeeTypeAlignment(ParamType).getAsAlign();
+        Attrs.addAlignmentAttr(Alignment);
+      }
+    }
+
     switch (FI.getExtParameterInfo(ArgNo).getABI()) {
     case ParameterABI::Ordinary:
       break;
 
     case ParameterABI::SwiftIndirectResult: {
       // Add 'sret' if we haven't already used it for something, but
       // only if the result is void.
       if (!hasUsedSRet && RetTy->isVoidType()) {

clang/test/CodeGenOpenCL/amdgpu-call-kernel.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
-// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture noundef writeonly %out)
+// CHECK: define{{.*}} amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture noundef writeonly align 4 %out)
 // CHECK: store i32 4, i32 addrspace(1)* %out, align 4
 
 kernel void test_kernel(global int *out)
 {
   out[0] = 4;
 }
 
 __kernel void test_call_kernel(__global int *out)
 {
   test_kernel(out);
 }

clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl

 // Emits block literal [[BL_GLOBAL]], invoke function [[INV_G]] and global block variable @block_G
 // COMMON: define internal spir_func void [[INV_G]](i8 addrspace(4)* %{{.*}}, i8 addrspace(3)* %{{.*}})
 const bl_t block_G = (bl_t) ^ (local void *a) {};
 
 void callee(int id, __global int *out) {
   out[id] = id;
 }
 
-// COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(i32 addrspace(1)* %{{.*}}, i32 addrspace(1)* %b, i32 %i)
+// COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(i32 addrspace(1)* align 4 %{{.*}}, i32 addrspace(1)* align 4 %b, i32 %i)
 kernel void device_side_enqueue(global int *a, global int *b, int i) {
   // COMMON: %default_queue = alloca %opencl.queue_t*
   queue_t default_queue;
   // COMMON: %flags = alloca i32
   unsigned flags = 0;
   // COMMON: %ndrange = alloca %struct.ndrange_t
   ndrange_t ndrange;
   // COMMON: %clk_event = alloca %opencl.clk_event_t*

clang/test/CodeGenOpenCL/kernel-param-alignment.cl

+// RUN: %clang_cc1 %s -cl-std=CL1.2 -emit-llvm -triple x86_64-unknown-unknown -o - | FileCheck %s
+
+// Test that pointer arguments to kernels are assumed to be ABI aligned.
+
+struct __attribute__((packed, aligned(1))) packed {
+  int i32;
+};
+
+typedef __attribute__((ext_vector_type(4))) int int4;
+typedef __attribute__((ext_vector_type(2))) float float2;
+
+kernel void test(
+    global int *i32,
+    global long *i64,
+    global int4 *v4i32,
+    global float2 *v2f32,
+    global void *v,
+    global struct packed *p) {
+// CHECK-LABEL: spir_kernel void @test(
+// CHECK-SAME: i32* nocapture noundef align 4 %i32,
+// CHECK-SAME: i64* nocapture noundef align 8 %i64,
+// CHECK-SAME: <4 x i32>* nocapture noundef align 16 %v4i32,
+// CHECK-SAME: <2 x float>* nocapture noundef align 8 %v2f32,
+// CHECK-SAME: i8* nocapture noundef %v,
+// CHECK-SAME: %struct.packed* nocapture noundef align 1 %p)
+}

clang/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl

typedef struct test_struct {
     short elementG;
     double elementH;
 } test_struct;
 
 kernel void test_single(int_single input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_single
 // CHECK: struct.int_single* nocapture {{.*}} byval(%struct.int_single)
-// CHECK: i32* nocapture noundef writeonly %output
+// CHECK: i32* nocapture noundef writeonly align 4 %output
  output[0] = input.a;
 }
 
 kernel void test_pair(int_pair input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_pair
 // CHECK: struct.int_pair* nocapture {{.*}} byval(%struct.int_pair)
-// CHECK: i32* nocapture noundef writeonly %output
+// CHECK: i32* nocapture noundef writeonly align 4 %output
  output[0] = (int)input.a;
  output[1] = (int)input.b;
 }
 
 kernel void test_kernel(test_struct input, global int* output) {
 // CHECK: spir_kernel
 // AMDGCN: define{{.*}} amdgpu_kernel void @test_kernel
 // CHECK: struct.test_struct* nocapture {{.*}} byval(%struct.test_struct)
-// CHECK: i32* nocapture noundef writeonly %output
+// CHECK: i32* nocapture noundef writeonly align 4 %output
  output[0] = input.elementA;
  output[1] = input.elementB;
  output[2] = (int)input.elementC;
  output[3] = (int)input.elementD;
  output[4] = (int)input.elementE;
  output[5] = (int)input.elementF;
  output[6] = (int)input.elementG;
  output[7] = (int)input.elementH;

clang/test/CodeGenOpenCL/spir-calling-conv.cl

 // RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s
 
 int get_dummy_id(int D);
 
 kernel void bar(global int *A);
 
 kernel void foo(global int *A)
-// CHECK: define{{.*}} spir_kernel void @foo(i32 addrspace(1)* noundef %A)
+// CHECK: define{{.*}} spir_kernel void @foo(i32 addrspace(1)* noundef align 4 %A)
 {
   int id = get_dummy_id(0);
   // CHECK: %{{[a-z0-9_]+}} = tail call spir_func i32 @get_dummy_id(i32 noundef 0)
   A[id] = id;
   bar(A);
-  // CHECK: tail call spir_kernel void @bar(i32 addrspace(1)* noundef %A)
+  // CHECK: tail call spir_kernel void @bar(i32 addrspace(1)* noundef align 4 %A)
 }
 
 // CHECK: declare spir_func i32 @get_dummy_id(i32 noundef)
-// CHECK: declare spir_kernel void @bar(i32 addrspace(1)* noundef)
+// CHECK: declare spir_kernel void @bar(i32 addrspace(1)* noundef align 4)