diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 5f8642df4c2a4..24dd8368a3cf0 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1280,15 +1280,17 @@ def SYCLRegisterNum : InheritableAttr { let Documentation = [SYCLRegisterNumDocs]; } -// Used by FE to mark ESIMD kernel pointer parameters which correspond to the +// Used by FE to mark SYCL kernel pointer parameters which correspond to the // original lambda's captured accessors. FE turns the attribute to some metadata -// required by the ESIMD Back-End. -// Not supposed to be used directly in the source - SYCL device compiler FE -// automatically adds it for ESIMD kernels, hence undocumented. -def SYCLSimdAccessorPtr : InheritableAttr { - // No spelling, as this attribute can't be created in the source code. +// required by the device back-end. +// This attribute does not require custom semantic handling +// hence we set the SemaHandler field to 0. +// The attribute is not for public consumption, and is an implicitly-created attribute +// that has no visible spelling, hence undocumented. +def SYCLAccessorPtr : Attr { + // This attribute has no spellings as it is only ever created implicitly. let Spellings = []; - let Subjects = SubjectList<[ParmVar]>; + let SemaHandler = 0; let Documentation = [Undocumented]; } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index e049069354838..e2b29efce106c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1728,10 +1728,16 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, // MDNode for the intel_buffer_location attribute. SmallVector argSYCLBufferLocationAttr; + // MDNode for listing SYCL kernel pointer arguments originating from + // accessors. + SmallVector argSYCLKernelRuntimeAligned; + // MDNode for listing ESIMD kernel pointer arguments originating from - // accessors + // accessors. SmallVector argESIMDAccPtrs; + bool isKernelArgAnAccessor = false; + if (FD && CGF) for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) { const ParmVarDecl *parm = FD->getParamDecl(i); @@ -1835,17 +1841,38 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SYCLBufferLocationAttr->getLocationID())) : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); + // If a kernel pointer argument comes from an accessor, we generate + // a new metadata(kernel_arg_runtime_aligned) to the kernel to indicate + // that this pointer has runtime allocated alignment. The value of any + // "kernel_arg_runtime_aligned" metadata element is 'true' for any kernel + // arguments that corresponds to the base pointer of an accessor and + // 'false' otherwise. + if (parm->hasAttr()) { + isKernelArgAnAccessor = true; + argSYCLKernelRuntimeAligned.push_back( + llvm::ConstantAsMetadata::get(CGF->Builder.getTrue())); + } else { + argSYCLKernelRuntimeAligned.push_back( + llvm::ConstantAsMetadata::get(CGF->Builder.getFalse())); + } + if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( - CGF->Builder.getInt1(parm->hasAttr()))); + CGF->Builder.getInt1(parm->hasAttr()))); } bool IsEsimdFunction = FD && FD->hasAttr(); - if (LangOpts.SYCLIsDevice && !IsEsimdFunction) + if (LangOpts.SYCLIsDevice && !IsEsimdFunction) { Fn->setMetadata("kernel_arg_buffer_location", llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr)); - else { + // Generate this metadata only if atleast one kernel argument is an + // accessor. + if (isKernelArgAnAccessor) + Fn->setMetadata( + "kernel_arg_runtime_aligned", + llvm::MDNode::get(VMContext, argSYCLKernelRuntimeAligned)); + } else { Fn->setMetadata("kernel_arg_addr_space", llvm::MDNode::get(VMContext, addressQuals)); Fn->setMetadata("kernel_arg_access_qual", diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ec733c8eccca2..434669019c51e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1962,10 +1962,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Additional processing is required for accessor type. void handleAccessorType(const CXXRecordDecl *RecordDecl, SourceLocation Loc) { handleAccessorPropertyList(Params.back(), RecordDecl, Loc); - if (KernelDecl->hasAttr()) - // In ESIMD, the kernels accessor's pointer argument needs to be marked. - Params.back()->addAttr( - SYCLSimdAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext())); // Get access mode of accessor. const auto *AccessorSpecializationDecl = cast(RecordDecl); @@ -1977,6 +1973,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { if (isReadOnlyAccessor(AccessModeArg)) Params.back()->addAttr( SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext())); + Params.back()->addAttr( + SYCLAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext())); } // All special SYCL objects must have __init method. We extract types for diff --git a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp new file mode 100644 index 0000000000000..500b5fac500a1 --- /dev/null +++ b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp @@ -0,0 +1,153 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s + +// This test checks if the metadata "kernel-arg-runtime-aligned" +// is generated if the kernel captures an accessor. + +#include "sycl.hpp" + +using namespace cl::sycl; + +queue q; + +int main() { + + using Accessor = + accessor; + Accessor acc[2]; + + accessor readOnlyAccessor; + + accessor + acc3; + + // kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, + // int*, sycl::range<1>, sycl::range<1>,sycl::id<1>. + q.submit([&](handler &h) { + h.single_task([=]() { + acc[1].use(); + }); + }); + + // kernel_readOnlyAcc parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>. + q.submit([&](handler &h) { + h.single_task([=]() { + readOnlyAccessor.use(); + }); + }); + + // kernel_B parameters : none. + q.submit([&](handler &h) { + h.single_task([=]() { + int result = 5; + }); + }); + + int a = 10; + + // kernel_C parameters : int. + q.submit([&](handler &h) { + h.single_task([=]() { + int x = a; + }); + }); + + // Using raw pointers to represent USM pointers. + // kernel_arg_runtime_aligned is not generated for raw pointers. + int *x; + float *y; + q.submit([&](handler &h) { + h.single_task([=]() { + *x = 42; + *y = 3.14; + }); + }); + + // Using local accessor as a kernel parameter. + // kernel_arg_runtime_aligned is generated for pointers from local accessors. + q.submit([&](handler &h) { + h.single_task([=]() { + acc3.use(); + }); + }); + + // kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*. + int *rawPtr; + q.submit([&](handler &h) { + h.single_task([=]() { + readOnlyAccessor.use(); + *rawPtr = 10; + }); + }); + + // Check if kernel_arg_accessor_ptr metadata is generated for ESIMD kernels that capture + // an accessor. + q.submit([&](handler &h) { + h.single_task([=]() __attribute__((sycl_explicit_simd)) { + readOnlyAccessor.use(); + }); + }); +} + +// Check kernel_A parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]], +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]]) +// CHECK-SAME: !kernel_arg_runtime_aligned !5 + +// Check kernel_readOnlyAcc parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_readOnlyAcc +// CHECK-SAME: i32 addrspace(1)* readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]] +// CHECK-SAME: !kernel_arg_runtime_aligned !14 + +// Check kernel_B parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_B +// CHECK-NOT: kernel_arg_runtime_aligned + +// Check kernel_C parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C +// CHECK-SAME: i32 [[MEM_ARG1:%[a-zA-Z0-9_]+]] +// CHECK-NOT: kernel_arg_runtime_aligned + +// Check usm_ptr parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}usm_ptr +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: float addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]] +// CHECK-NOT: kernel_arg_runtime_aligned + +// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor +// CHECK-SAME: float addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.cl::sycl::range.5"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range.5"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id.6"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]] +// CHECK-SAME: !kernel_arg_runtime_aligned !14 + +// Check kernel_acc_raw_ptr parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr +// CHECK-SAME: i32 addrspace(1)* readonly [[MEM_ARG1:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]] +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]] +// CHECK-SAME: !kernel_arg_runtime_aligned !26 + +// Check esimd_kernel_with_acc parameters +// CHECK: define {{.*}}spir_kernel void @{{.*}}esimd_kernel_with_acc +// CHECK-SAME: !kernel_arg_accessor_ptr + +// Check kernel-arg-runtime-aligned metadata. +// The value of any metadata element is 1 for any kernel arguments +// that corresponds to the base pointer of an accessor and 0 otherwise. +// CHECK: !5 = !{i1 true, i1 false, i1 false, i1 false, i1 true, i1 false, i1 false, i1 false} +// CHECK: !14 = !{i1 true, i1 false, i1 false, i1 false} +// CHECK: !26 = !{i1 true, i1 false, i1 false, i1 false, i1 false} diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 6f9cc4091fc4a..1f88f52c6fbcf 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -114,10 +114,12 @@ int main() { // Check Kernel_Accessor parameters // CHECK: FunctionDecl {{.*}}Kernel_Accessor{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *' +// CHECK-NEXT: SYCLAccessorPtrAttr // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::id<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *' +// CHECK-NEXT: SYCLAccessorPtrAttr // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'sycl::id<1>' @@ -165,10 +167,12 @@ int main() { // Check Kernel_StructAccArray parameters // CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *' +// CHECK-NEXT: SYCLAccessorPtrAttr // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::id<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *' +// CHECK-NEXT: SYCLAccessorPtrAttr // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::id<1>'