Skip to content
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

[SYCL] Avoid alignment on kernel pointer parameters #11979

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

[SYCL] Avoid alignment on kernel pointer parameters #11979

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
6 changes: 4 additions & 2 deletions clang/lib/CodeGen/CGCall.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2850,8 +2850,10 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
// > 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()) {
//
// Don't do this for SYCL, as this assumption does not hold.
if (!getLangOpts().SYCLIsDevice && TargetDecl &&
TargetDecl->hasAttr<OpenCLKernelAttr>() && ParamType->isPointerType()) {
Comment on lines +2853 to +2856
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@premanandrao, do I get it right that alignment is not guaranteed for USM allocations only?
Is it true for L0 only or OpenCL is impacted as well?

I'm surprised to see the deviation from OpenCL properties. It might hard to justify in upstream. If SYCL compiler doesn't genuine OpenCL kernel, can we continue using OpenCLKernelAttr or better to have a SYCL specific attribute?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is not guaranteed for USM allocations, but since we don't do pointer analysis, we can't deduce the alignment in general.

Good questions about L0 vs OpenCL. I agree with your suggestion that perhaps we use a different SYCL attribute if it applies to OpenCL.

I have added @GarveyJoe and @ajaykumarkannan to the PR; they had identified and requested this change. I would like to have their thoughts on this too.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

According to https://registry.khronos.org/OpenCL/extensions/intel/cl_intel_unified_shared_memory.html. USM allocation alignment requirements match OpenCL buffer (i.e. it must be a power of two and must be equal to or smaller than the size of the largest data type supported by any OpenCL device in context), so we can re-use OpenCL kernel logic as-is for pointers to USM allocations. We can argue about whether OpenCL logic is correct, but I don't think it should cause the difference between OpenCL and SYCL.

I don't see similar alignment requirements for Level Zero though. Level Zero spec only requires alignment value to be a power of two. @bashbaug, do you know if Level Zero memory allocation functions have additional alignment guarantees like OpenCL?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@bader, the wording you're looking at regarding alignment is about the alignment of the pointer returned by the allocation functions such as clSharedMemAllocINTEL. There is no requirement that the kernel argument passed in via clSetKernelArgMemPointerINTEL has that same alignment. The only restriction on the pointer passed to clSetKernelArgMemPointerINTEL is that it is somewhere within an allocation returned by one of the allocation functions:

Otherwise, the pointer value must be NULL or must point into a Unified Shared Memory allocation returned by clHostMemAllocINTEL, clDeviceMemAllocINTEL, or clSharedMemAllocINTEL.

As a result, the following code is legal OpenCL:

// Kernel
kernel void foo(global int *a) {
  global char *b = (char *)(a);
  *b =  ...;
}

// Host
char *p = (char *) clHostMemAllocINTEL(context, NULL, 2, 0, NULL);
p = &(p[1]);
clSetKernelArgMemPointerINTEL(kernel, 0, p);

And certainly in this case the kernel argument will not have alignment any higher than that of a char.

@premanandrao, since this same problem can be exposed in OpenCL, as my example demonstrates, I don't think your code should be SYCL-specific.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Note, there is a line in the OpenCL C spec saying:

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.

A similar line also exists in the OpenCL SPIR-V environment spec:

For OpTypePointer arguments to a function, the compiler may assume that the pointer is appropriately aligned as required by the Type that the pointer points to.

The example above still might be OK, but because the int* kernel argument is not aligned to sizeof(int) == 4 bytes things could easily go wrong.

Does SYCL (or C++, generally) make similar guarantees?

do you know if Level Zero memory allocation functions have additional alignment guarantees like OpenCL?

The Level Zero memory allocation functions have a similar alignment parameter as the OpenCL allocation functions. The Level Zero spec doesn't seem to explicitly say what the behavior is when passing zero as the alignment, but I'm 99% sure it behaves the same as OpenCL, by choosing an implementation-defined alignment that is big enough for all basic data types.

Copy link
Contributor

@GarveyJoe GarveyJoe Dec 7, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

After much digging through the spec, I've concluded that this optimization is actually legal in all C++ programs (and thus in SYCL as well). Without ever saying so explicitly, the language standard goes to great lengths to ensure that any pointer that doesn't have at least the alignment of the type it points to has either an undefined value or produces undefined behaviour even if the pointer is never dereferenced. However, it seems that this is not the approach that clang has taken. This LLVM mailing list discussion started by John McCall in 2016 seems to summarize clang's current position: https://groups.google.com/g/llvm-dev/c/eJRto1ipCYQ. In that thread he proposes that clang maintain a more relaxed position than the C++ standard: that it is only UB to dereference an unaligned pointer. I can't find anything formal in the clang docs to indicate that his proposal was accepted but even in present day clang does not emit the alignment attributes that the stricter definition from the C++ standard would allow. It instead only emits alignment at access sites. At the very least it seems the clang community has tacitly accepted John's proposal.

Now we have to decide if we want to take advantage of the stronger guarantees of the standard or follow clang's looser direction. I suspect we might get push back from the community if we try to upstream code that takes advantage of this guarantee.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The SYCL 2020 spec, in section 5.5 (Built-in scalar data types), requires scalar fundamental data types to have the same size and alignment for the host and device. The alignment annotations look correct to me as is.

I believe the stronger guarantee exists to allow for an implementation to diagnose the creation of an invalid pointer as opposed to having to wait until the pointer is dereferenced.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Following further offline discussion, I now agree with Joe that this is a good change. Without it, code might behave differently in a kernel than in another device function and that just seems weird and unnecessary. I don't think the optimization opportunity is significant.

One of the things that was helpful for me in reaching this conclusion is that alias annotations in LLVM IR are coalesced; when there are multiple relevant annotations (e.g., on a parameter with a pointer type and on a load/store that uses that pointer), code gen can use the more strict one.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think the optimization opportunity is significant.

One of the things that was helpful for me in reaching this conclusion is that alias annotations in LLVM IR are coalesced; when there are multiple relevant annotations (e.g., on a parameter with a pointer type and on a load/store that uses that pointer), code gen can use the more strict one.

I strongly recommend testing this claim using available means.

@jingwan2, FYI.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Testing is always a good idea :)

The reason I think the optimization opportunity is not significant is because the alias information is (currently) lost as soon as one of these pointers is passed to another function (though subject to inlining considerations I'm sure). At any rate, it would be good to get input from someone with actual optimization experience.

Comment on lines +2854 to +2856
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Assuming that the OpenCLKernelAttr attribute is only used for OpenCL and SYCL, perhaps it makes sense to restrict the conditional to matching OpenCL specifically rather than any language extension other than SYCL.

Suggested change
// Don't do this for SYCL, as this assumption does not hold.
if (!getLangOpts().SYCLIsDevice && TargetDecl &&
TargetDecl->hasAttr<OpenCLKernelAttr>() && ParamType->isPointerType()) {
if (getLangOpts().OpenCL && TargetDecl &&
TargetDecl->hasAttr<OpenCLKernelAttr>() && ParamType->isPointerType()) {

QualType PTy = ParamType->getPointeeType();
if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {
llvm::Align Alignment =
Expand Down
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
Expand Up @@ -670,5 +670,8 @@ class image {
}
};

template <typename T> T *malloc_shared(unsigned long size);
void free(void *ptr);

} // namespace _V1
} // namespace sycl
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/accessor-readonly.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@ void f0(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1>
// CHECK: spir_kernel{{.*}}f1_kernel
// CHECK-NOT: readonly
// CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
// CHECK-SAME: readonly align 4 %_arg_read_acc
// CHECK-SAME: readonly %_arg_read_acc
void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](sycl::handler &cgh) {
auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
Expand All @@ -25,7 +25,7 @@ void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1>
}

// CHECK: spir_kernel{{.*}}f2_kernel
// CHECK-SAME: readonly align 4 %_arg_read_acc
// CHECK-SAME: readonly %_arg_read_acc
// CHECK-NOT: readonly
// CHECK-SAME: %_arg_write_acc
void f2(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,11 +26,11 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
// CHECK-SAME: i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]],
// CHECK-SAME: i32 noundef [[ARG_B:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC1_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly [[ACC1_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly align 1 [[ACC2_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly [[ACC2_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@ int main() {
}

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/intel-restrict.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,18 +11,18 @@ int main() {
int *c;
kernel<class kernel_restrict>(
[ a, b, c ]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0]; });
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict(ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}})
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict(ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}})

int *d;
int *e;
int *f;

kernel<class kernel_norestrict>(
[d, e, f]() { f[0] = d[0] + e[0]; });
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_norestrict(ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}}, ptr addrspace(1) noundef align 4 %{{.*}})
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_norestrict(ptr addrspace(1) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}})

int g = 42;
kernel<class kernel_restrict_other_types>(
[ a, b, c, g ]() [[intel::kernel_args_restrict]] { c[0] = a[0] + b[0] + g; });
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict_other_types(ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, ptr addrspace(1) noalias noundef align 4 %{{.*}}, i32 noundef %{{.*}})
// CHECK: define {{.*}}spir_kernel {{.*}}kernel_restrict_other_types(ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, ptr addrspace(1) noalias noundef %{{.*}}, i32 noundef %{{.*}})
}
18 changes: 9 additions & 9 deletions clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -102,11 +102,11 @@ int main() {

// Check kernel_A parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
Expand All @@ -115,7 +115,7 @@ int main() {

// Check kernel_readOnlyAcc parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_readOnlyAcc
// CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
Expand All @@ -135,21 +135,21 @@ int main() {

// Check usm_ptr parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}usm_ptr
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]]
// CHECK-NOT: kernel_arg_runtime_aligned
// CHECK-NOT: kernel_arg_exclusive_ptr

// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]]
// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]

// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
// CHECK-SAME: ptr addrspace(3) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(3) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
Expand All @@ -158,11 +158,11 @@ int main() {

// Check kernel_acc_raw_ptr parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
// CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]]
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD3:]]
// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD3]]

Expand Down
52 changes: 52 additions & 0 deletions clang/test/CodeGenSYCL/kernel-arg-align.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,52 @@
// RUN: %clang_cc1 -fsycl-is-device -O0 -internal-isystem %S/Inputs -triple spir64 -emit-llvm -o - %s | FileCheck %s
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
// RUN: %clang_cc1 -fsycl-is-device -O0 -internal-isystem %S/Inputs -triple spir64 -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -fsycl-is-device -Xclang -disable-llvm-passes -internal-isystem %S/Inputs -triple spir64 -emit-llvm -o - %s | FileCheck %s


// Test that the pointer parameters generated for the kernel do not
// have alignment on them.

#include "sycl.hpp"

using namespace sycl;

struct S;

void Test() {
struct MyIP {
char* a;
int* b;
double* c;

void operator()() const {
*((int *) a) = 1; // 1 on arg, 4 on site
*((double *) b) = 2; // 4 on arg, 8 on site
*((char *) c) = 3; // 8 on arg, 1 on site
}
};

constexpr int kN = 8;
auto host_array_A =
malloc_shared<char>(kN);

auto host_array_B =
malloc_shared<int>(kN);

auto host_array_C =
malloc_shared<double>(kN);

for (int i = 0; i < kN; i++) {
host_array_A[i] = i;
host_array_B[i] = i * 2;
}

sycl::kernel_single_task<S>(MyIP{host_array_A, host_array_B, host_array_C});

free(host_array_A);
free(host_array_B);
free(host_array_C);
}

int main() {
Test();
return 0;
}

// CHECK: define {{.*}} spir_kernel void @_ZTS1S(ptr addrspace(1) noundef %_arg_a, ptr addrspace(1) noundef %_arg_b, ptr addrspace(1) noundef %_arg_c)
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ void test(int val) {
}

// ALL: define dso_local{{ spir_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
// ALL-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
// ALL-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef %_arg__specialization_constants_buffer)
// ALL: %kh = alloca %"class.sycl::_V1::kernel_handler", align 1

// NONATIVESUPPORT: %[[KH:[0-9]+]] = load ptr addrspace(1), ptr %_arg__specialization_constants_buffer.addr, align 8
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,11 +25,11 @@ int main() {

// Check kernel_A parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,11 +28,11 @@ int main() {

// CHECK kernel_C parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr addrspace(1) noundef [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/stream.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@
// CHECK: %[[ID_TYPE:"struct.*sycl::_V1::id"]]

// CHECK: define dso_local spir_kernel void @{{.*}}StreamTester
// CHECK-SAME: ptr addrspace(1) noundef align 1 [[ACC_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr addrspace(1) noundef [[ACC_DATA:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+]],
// CHECK-SAME: ptr noundef byval(%[[ID_TYPE]]) align 4 [[ACC_ID:%[a-zA-Z0-9_]+]],
Expand Down