From d8ad40296bfb677c90f83d43ea104993a9ddb25a Mon Sep 17 00:00:00 2001 From: y Date: Mon, 11 May 2026 15:47:00 -0700 Subject: [PATCH 1/3] [SYCL] Preserve restrict on free function kernel pointer args Preserve top-level pointer qualifiers when rewriting SYCL free function kernel pointer parameters into global address space, so restrict lowers to LLVM noalias on generated kernel arguments. Add device-code coverage for restrict-qualified free function kernel and helper parameters. --- clang/lib/Sema/SemaSYCL.cpp | 4 ++- .../free_function_kernels_restrict.cpp | 30 +++++++++++++++++++ 2 files changed, 33 insertions(+), 1 deletion(-) create mode 100644 sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9be0172eb7971..704a9b5913a85 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2477,6 +2477,7 @@ static QualType ModifyAddressSpace(SemaSYCL &SemaSYCLRef, QualType Ty) { // same type as field but global address space, because OpenCL requires it. // Function pointers should have program address space. This is set in // CodeGen. + Qualifiers PtrQuals = Ty.getQualifiers(); QualType PointeeTy = Ty->getPointeeType(); Qualifiers Quals = PointeeTy.getQualifiers(); LangAS AS = Quals.getAddressSpace(); @@ -2487,7 +2488,8 @@ static QualType ModifyAddressSpace(SemaSYCL &SemaSYCLRef, QualType Ty) { Quals.setAddressSpace(LangAS::sycl_global); PointeeTy = SemaSYCLRef.getASTContext().getQualifiedType( PointeeTy.getUnqualifiedType(), Quals); - return SemaSYCLRef.getASTContext().getPointerType(PointeeTy); + QualType PtrTy = SemaSYCLRef.getASTContext().getPointerType(PointeeTy); + return SemaSYCLRef.getASTContext().getQualifiedType(PtrTy, PtrQuals); } // This visitor is used to traverse a non-decomposed record/array to diff --git a/sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp b/sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp new file mode 100644 index 0000000000000..c3396e982974a --- /dev/null +++ b/sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp @@ -0,0 +1,30 @@ +// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +// CHECK-IR-DAG: define dso_local spir_func noundef i32 @_Z15helper_restrictPKi(ptr addrspace(4) noalias +SYCL_EXTERNAL int helper_restrict(const int *__restrict__ input) { + return input[0]; +} + +// CHECK-IR-DAG: define dso_local spir_func void @_Z25helper_const_ptr_restrictPiPKi(ptr addrspace(4) noalias{{[^,]*}}, ptr addrspace(4) noalias +SYCL_EXTERNAL void +helper_const_ptr_restrict(int *const __restrict__ output, + const int *const __restrict__ input) { + output[0] = input[0]; +} + +// CHECK-IR: define dso_local spir_kernel void @_Z43__sycl_kernel_free_function_kernel_restrictPiPKiS_S_S1_PKv(ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) {{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void free_function_kernel_restrict(int *__restrict__ output, + const int *__restrict__ input, int *plain, + int *const __restrict__ output_const, + const int *const __restrict__ input_const, + const void *const __restrict__ opaque) { + output_const[0] = plain[0]; + helper_const_ptr_restrict(output, input_const); + output[0] += helper_restrict(input) + output_const[0] + (opaque != nullptr); +} From b16b33480b5b7ddc96716da8505a3937096083d1 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 12 May 2026 06:22:36 -0700 Subject: [PATCH 2/3] Moved test under appropriate directory --- .../test/CodeGenSYCL}/free_function_kernels_restrict.cpp | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename {sycl/test/check_device_code/extensions => clang/test/CodeGenSYCL}/free_function_kernels_restrict.cpp (100%) diff --git a/sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp b/clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp similarity index 100% rename from sycl/test/check_device_code/extensions/free_function_kernels_restrict.cpp rename to clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp From ae9dafe70b761a98a9931b3e7ef0b71beef5c90d Mon Sep 17 00:00:00 2001 From: y Date: Tue, 12 May 2026 15:53:54 -0700 Subject: [PATCH 3/3] Fix comment --- .../CodeGenSYCL/free_function_kernels_restrict.cpp | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp b/clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp index c3396e982974a..ff5da868961d3 100644 --- a/clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp +++ b/clang/test/CodeGenSYCL/free_function_kernels_restrict.cpp @@ -1,24 +1,21 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2020 -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -#include -#include - -namespace syclexp = sycl::ext::oneapi::experimental; +#include "sycl.hpp" // CHECK-IR-DAG: define dso_local spir_func noundef i32 @_Z15helper_restrictPKi(ptr addrspace(4) noalias -SYCL_EXTERNAL int helper_restrict(const int *__restrict__ input) { +__attribute__((sycl_device)) int helper_restrict(const int *__restrict__ input) { return input[0]; } // CHECK-IR-DAG: define dso_local spir_func void @_Z25helper_const_ptr_restrictPiPKi(ptr addrspace(4) noalias{{[^,]*}}, ptr addrspace(4) noalias -SYCL_EXTERNAL void +__attribute__((sycl_device)) void helper_const_ptr_restrict(int *const __restrict__ output, const int *const __restrict__ input) { output[0] = input[0]; } // CHECK-IR: define dso_local spir_kernel void @_Z43__sycl_kernel_free_function_kernel_restrictPiPKiS_S_S1_PKv(ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) {{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias{{[^,]*}}, ptr addrspace(1) noalias -SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]] void free_function_kernel_restrict(int *__restrict__ output, const int *__restrict__ input, int *plain, int *const __restrict__ output_const,