diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 87a0e04c88466..c4fa48ac729f2 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -296,6 +296,7 @@ ENUM_LANGOPT(SYCLRangeRounding, SYCLRangeRoundingPreference, 2, LANGOPT(SYCLExperimentalRangeRounding, 1, 0, NotCompatible, "Use experimental parallel for range rounding") LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, NotCompatible, "Enable diagnostics that require the SYCL integration header") LANGOPT(SYCLRTCMode, 1, 0, NotCompatible, "Compile in RTC mode") +LANGOPT(SYCLForceGlobalASInKernelArgs, 1, 0, NotCompatible, "Force global address space in SYCL kernel arguments") LANGOPT(HIPUseNewLaunchAPI, 1, 0, NotCompatible, "Use new kernel launching API for HIP") LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)") diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index 6a67c0e7af06f..2c724364fcb39 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -9749,6 +9749,9 @@ defm offload_use_alloca_addrspace_for_srets : BoolFOption<"offload-use-alloca-ad DefaultTrue, PosFlag, NegFlag>; +def fsycl_force_global_as_in_kernel_args : Flag<["-"], "fsycl-force-global-as-in-kernel-args">, + HelpText<"Force global address space for USM pointers within SYCL kernel arguments">, + MarshallingInfoFlag>; def fsycl_remangle_libspirv : Flag<["-"], "fsycl-remangle-libspirv">, diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 2609dc59fc3ac..7e8dcdac8f2f4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2333,7 +2333,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { SemaSYCLRef.getASTContext())); PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - if (!RD->hasAttr()) + if (!RD->hasAttr() && + SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs) RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit( SemaSYCLRef.getASTContext())); } @@ -2369,7 +2370,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { PointerStack.back() = true; - if (!RD->hasAttr()) + if (!RD->hasAttr() && + SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs) RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit( SemaSYCLRef.getASTContext())); } @@ -2388,7 +2390,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { PointerStack.back() = true; - if (!RD->hasAttr()) + if (!RD->hasAttr() && + SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs) RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit( SemaSYCLRef.getASTContext())); } @@ -2418,7 +2421,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { PointerStack.back() = true; - if (!RD->hasAttr()) + if (!RD->hasAttr() && + SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs) RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit( SemaSYCLRef.getASTContext())); } @@ -2452,7 +2456,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - if (!FD->hasAttr()) + if (!FD->hasAttr() && + SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs) FD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit( SemaSYCLRef.getASTContext())); PointerStack.back() = true; diff --git a/clang/test/CodeGenSYCL/array-kernel-param.cpp b/clang/test/CodeGenSYCL/array-kernel-param.cpp index 1948f4a3abf1f..0cee7422e339b 100644 --- a/clang/test/CodeGenSYCL/array-kernel-param.cpp +++ b/clang/test/CodeGenSYCL/array-kernel-param.cpp @@ -1,16 +1,17 @@ -// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s #include "sycl.hpp" +// CHECK: %struct.__wrapper_class = type { [2 x i32] } +// CHECK: %class.anon = type { [2 x i32] } +// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(4)] } +// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] } + sycl::queue myQueue; using namespace sycl; -// CHECK: %struct.__wrapper_class = type { [2 x i32] } -// CHECK: %class.anon = type { [2 x i32] } -// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(1)] } -// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] } int main() { int Array[2]; @@ -21,24 +22,6 @@ int main() { }); }); -// CHECK-LABEL: @{{.*}}IntArray(ptr {{.*}}byval(%struct.__wrapper_class) -// CHECK: %__SYCLKernel = alloca %class.anon, align 4 -// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4) -// CHECK: %_arg_Array.ascast = addrspacecast ptr %_arg_Array to ptr addrspace(4) -// CHECK: %Array = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0 -// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class, ptr addrspace(4) %_arg_Array.ascast, i32 0, i32 0 -// CHECK: %arrayinit.begin = getelementptr inbounds [2 x i32], ptr addrspace(4) %Array, i64 0, i64 0 -// CHECK: br label %arrayinit.body -// CHECK: arrayinit.body: ; preds = %arrayinit.body, %entry -// CHECK: %arrayinit.index = phi i64 [ 0, %entry ], [ %arrayinit.next, %arrayinit.body ] -// CHECK: %1 = getelementptr inbounds i32, ptr addrspace(4) %arrayinit.begin, i64 %arrayinit.index -// CHECK: %arrayidx = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) %0, i64 0, i64 %arrayinit.index -// CHECK: %2 = load i32, ptr addrspace(4) %arrayidx, align 4 -// CHECK: store i32 %2, ptr addrspace(4) %1, align 4 -// CHECK: %arrayinit.next = add nuw i64 %arrayinit.index, 1 -// CHECK: %arrayinit.done = icmp eq i64 %arrayinit.next, 2 -// CHECK: br i1 %arrayinit.done, label %arrayinit.end, label %arrayinit.body -// CHECK: arrayinit.end: ; preds = %arrayinit.body int *ArrayOfPointers[2]; myQueue.submit([&](sycl::handler &h) { @@ -47,20 +30,83 @@ int main() { int local = *ArrayOfPointers[1]; }); }); -// CHECK-LABEL: @{{.*}}PtrArray(ptr {{.*}}byval(%struct.__wrapper_class.0) -// CHECK: %__SYCLKernel = alloca %class.anon.1, align 8 -// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4) -// CHECK: %_arg_ArrayOfPointers.ascast = addrspacecast ptr %_arg_ArrayOfPointers to ptr addrspace(4) -// CHECK: %ArrayOfPointers = getelementptr inbounds nuw %class.anon.1, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0 -// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0 -// CHECK: %arrayidx = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %0, i64 0, i64 0 -// CHECK: %1 = load ptr addrspace(1), ptr addrspace(4) %arrayidx, align 8 -// CHECK: %2 = addrspacecast ptr addrspace(1) %1 to ptr addrspace(4) -// CHECK: store ptr addrspace(4) %2, ptr addrspace(4) %ArrayOfPointers, align 8 -// CHECK: %arrayinit.element = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) %ArrayOfPointers, i64 1 -// CHECK: %3 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0 -// CHECK: %arrayidx1 = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %3, i64 0, i64 1 -// CHECK: %4 = load ptr addrspace(1), ptr addrspace(4) %arrayidx1, align 8 -// CHECK: %5 = addrspacecast ptr addrspace(1) %4 to ptr addrspace(4) -// CHECK: store ptr addrspace(4) %5, ptr addrspace(4) %arrayinit.element, align 8 } +// CHECK-LABEL: define weak_odr spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8IntArray( +// CHECK-SAME: ptr noundef byval([[STRUCT___WRAPPER_CLASS:%.*]]) align 4 [[_ARG_ARRAY:%.*]]) #[[ATTR0:[0-9]+]] comdat !srcloc [[META12:![0-9]+]] !kernel_arg_buffer_location [[META13:![0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*]]: +// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4 +// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ARRAY_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ARRAY]] to ptr addrspace(4) +// CHECK-NEXT: [[ARRAY:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___WRAPPER_CLASS]], ptr addrspace(4) [[_ARG_ARRAY_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x i32], ptr addrspace(4) [[ARRAY]], i64 0, i64 0 +// CHECK-NEXT: br label %[[ARRAYINIT_BODY:.*]] +// CHECK: [[ARRAYINIT_BODY]]: +// CHECK-NEXT: [[ARRAYINIT_INDEX:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[ARRAYINIT_NEXT:%.*]], %[[ARRAYINIT_BODY]] ] +// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[ARRAYINIT_BEGIN]], i64 [[ARRAYINIT_INDEX]] +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) [[TMP0]], i64 0, i64 [[ARRAYINIT_INDEX]] +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4 +// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[TMP1]], align 4 +// CHECK-NEXT: [[ARRAYINIT_NEXT]] = add nuw i64 [[ARRAYINIT_INDEX]], 1 +// CHECK-NEXT: [[ARRAYINIT_DONE:%.*]] = icmp eq i64 [[ARRAYINIT_NEXT]], 2 +// CHECK-NEXT: br i1 [[ARRAYINIT_DONE]], label %[[ARRAYINIT_END:.*]], label %[[ARRAYINIT_BODY]] +// CHECK: [[ARRAYINIT_END]]: +// CHECK-NEXT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(8) [[__SYCLKERNEL_ASCAST]]) #[[ATTR2:[0-9]+]] +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv( +// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable_or_null(8) [[THIS:%.*]]) #[[ATTR1:[0-9]+]] align 2 !srcloc [[META12]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr addrspace(4), align 8 +// CHECK-NEXT: [[LOCAL:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[THIS_ADDR_ASCAST:%.*]] = addrspacecast ptr [[THIS_ADDR]] to ptr addrspace(4) +// CHECK-NEXT: [[LOCAL_ASCAST:%.*]] = addrspacecast ptr [[LOCAL]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(4) [[THIS]], ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[THIS1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[ARRAY:%.*]] = getelementptr inbounds nuw [[CLASS_ANON:%.*]], ptr addrspace(4) [[THIS1]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr addrspace(4) [[ARRAY]], i64 0, i64 1 +// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4 +// CHECK-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[LOCAL_ASCAST]], align 4 +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define weak_odr spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E8PtrArray( +// CHECK-SAME: ptr noundef byval([[STRUCT___WRAPPER_CLASS_0:%.*]]) align 8 [[_ARG_ARRAYOFPOINTERS:%.*]]) #[[ATTR0]] comdat !srcloc [[META14:![0-9]+]] !kernel_arg_buffer_location [[META13]] { +// CHECK-NEXT: [[ENTRY:.*]]: +// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 8 +// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ARRAYOFPOINTERS_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ARRAYOFPOINTERS]] to ptr addrspace(4) +// CHECK-NEXT: [[ARRAYOFPOINTERS:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___WRAPPER_CLASS_0]], ptr addrspace(4) [[_ARG_ARRAYOFPOINTERS_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[ARRAYOFPOINTERS]], i64 0, i64 0 +// CHECK-NEXT: br label %[[ARRAYINIT_BODY:.*]] +// CHECK: [[ARRAYINIT_BODY]]: +// CHECK-NEXT: [[ARRAYINIT_INDEX:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[ARRAYINIT_NEXT:%.*]], %[[ARRAYINIT_BODY]] ] +// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) [[ARRAYINIT_BEGIN]], i64 [[ARRAYINIT_INDEX]] +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x ptr addrspace(4)], ptr addrspace(4) [[TMP0]], i64 0, i64 [[ARRAYINIT_INDEX]] +// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8 +// CHECK-NEXT: store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[TMP1]], align 8 +// CHECK-NEXT: [[ARRAYINIT_NEXT]] = add nuw i64 [[ARRAYINIT_INDEX]], 1 +// CHECK-NEXT: [[ARRAYINIT_DONE:%.*]] = icmp eq i64 [[ARRAYINIT_NEXT]], 2 +// CHECK-NEXT: br i1 [[ARRAYINIT_DONE]], label %[[ARRAYINIT_END:.*]], label %[[ARRAYINIT_BODY]] +// CHECK: [[ARRAYINIT_END]]: +// CHECK-NEXT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[__SYCLKERNEL_ASCAST]]) #[[ATTR2]] +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv( +// CHECK-SAME: ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[THIS:%.*]]) #[[ATTR1]] align 2 !srcloc [[META14]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr addrspace(4), align 8 +// CHECK-NEXT: [[LOCAL:%.*]] = alloca i32, align 4 +// CHECK-NEXT: [[THIS_ADDR_ASCAST:%.*]] = addrspacecast ptr [[THIS_ADDR]] to ptr addrspace(4) +// CHECK-NEXT: [[LOCAL_ASCAST:%.*]] = addrspacecast ptr [[LOCAL]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(4) [[THIS]], ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[THIS1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[ARRAYOFPOINTERS:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1:%.*]], ptr addrspace(4) [[THIS1]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[ARRAYOFPOINTERS]], i64 0, i64 1 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[TMP0]], align 4 +// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[LOCAL_ASCAST]], align 4 +// CHECK-NEXT: ret void diff --git a/clang/test/CodeGenSYCL/free_function_kernel_params.cpp b/clang/test/CodeGenSYCL/free_function_kernel_params.cpp index 48ddd76ec99f3..d75ef7bf21123 100644 --- a/clang/test/CodeGenSYCL/free_function_kernel_params.cpp +++ b/clang/test/CodeGenSYCL/free_function_kernel_params.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \ -// RUN: -emit-llvm %s -o - | FileCheck %s +// RUN: -emit-llvm %s -o - | FileCheck --check-prefixes=CHECK,GEN-AS %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \ +// RUN: -fsycl-force-global-as-in-kernel-args -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,GLOB-AS // This test checks parameter IR generation for free functions with parameters // of non-decomposed struct type, work group memory type, dynamic work group memory type // and special types. @@ -47,16 +49,18 @@ void ff_6(KArgWithPtrArray KArg) { template void ff_6(KArgWithPtrArray KArg); -// CHECK: %struct.NoPointers = type { i32 } -// CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) } -// CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers } -// CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) } -// CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.4 } -// CHECK: %struct.__generated_Pointers.4 = type { ptr addrspace(1), ptr addrspace(1) } -// CHECK: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] } -// CHECK: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] } -// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3) -// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg) +// GEN-AS: %struct.NoPointers = type { i32 } +// GEN-AS: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) } +// GEN-AS: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers } +// GLOB-AS: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) } +// GLOB-AS: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.4 } +// GLOB-AS: %struct.__generated_Pointers.4 = type { ptr addrspace(1), ptr addrspace(1) } +// GLOB-AS: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] } +// GLOB-AS: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] } +// GLOB-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3) +// GLOB-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg) +// GEN-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.Agg) align 8 %__arg_S3) +// GEN-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.KArgWithPtrArray) align 8 %__arg_KArg) __attribute__((sycl_device)) [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] diff --git a/clang/test/CodeGenSYCL/generated-types-initialization.cpp b/clang/test/CodeGenSYCL/generated-types-initialization.cpp index b5e7f305a93ae..644e96ab33400 100644 --- a/clang/test/CodeGenSYCL/generated-types-initialization.cpp +++ b/clang/test/CodeGenSYCL/generated-types-initialization.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck %s // This test checks that compiler generates correct code when kernel arguments // are structs that contain pointers but not decomposed. diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index fc9e0957a7010..9ab6da7d50b20 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck --check-prefixes=CHECK,GEN-AS %s +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck --check-prefix=CHECK,GLOB-AS %s #include "Inputs/sycl.hpp" @@ -44,14 +45,16 @@ int main() { // CHECK: %struct.base = type { i32, %class.InnerField } // CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } // CHECK: %class.InnerFieldBase = type { i32 } -// CHECK: %class.__generated_second_base = type { ptr addrspace(1), [2 x ptr addrspace(1)] } -// CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> +// GLOB-AS: %class.__generated_second_base = type { ptr addrspace(1), [2 x ptr addrspace(1)] } +// GLOB-AS: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> // CHECK: %class.second_base = type { ptr addrspace(4), [2 x ptr addrspace(4)] } +// GEN-AS: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> // Check kernel paramters // CHECK: define {{.*}}spir_kernel void @{{.*}}derived // CHECK-SAME: ptr noundef byval(%struct.base) align 4 %_arg__base -// CHECK-SAME: ptr noundef byval(%class.__generated_second_base) align 8 %_arg__base1 +// GEN-AS-SAME: ptr noundef byval(%class.second_base) align 8 %_arg__base1 +// GLOB-AS-SAME: ptr noundef byval(%class.__generated_second_base) align 8 %_arg__base1 // CHECK-SAME: i32 noundef %_arg_a // Check allocas for kernel parameters and local functor object diff --git a/clang/test/CodeGenSYCL/pointers-in-structs.cpp b/clang/test/CodeGenSYCL/pointers-in-structs.cpp index 690ce34e8980e..74e2e24add89e 100644 --- a/clang/test/CodeGenSYCL/pointers-in-structs.cpp +++ b/clang/test/CodeGenSYCL/pointers-in-structs.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck %s // This test checks that compiler generates correct address spaces for pointer // kernel arguments that are wrapped by struct. Generated class should retain diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index ae6ceba69ccfc..c3c132ac85322 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck --check-prefixes=CHECK,GEN-AS %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -fsycl-force-global-as-in-kernel-args -sycl-std=2020 %s | FileCheck --check-prefixes=CHECK,GLOB-AS %s // This test checks that compiler generates correct kernel arguments for // arrays, Accessor arrays, and structs containing Accessors. @@ -157,23 +158,40 @@ int main() { // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' + +// Array initilializer if no address space modification +// GEN-AS-NEXT: ArrayInitLoopExpr {{.*}} 'int *[2]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' '__wrapper_class' +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'int *' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int **' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' '__wrapper_class' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' + +// GLOB-AS-NEXT: InitListExpr {{.*}} 'int *[2]' + +// Otherwise per-element initialization for address space modifications. // Initializer for ArrayOfPointers[0] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 + // Initializer for ArrayOfPointers[1] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 // 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>) __attribute__((device_kernel))' @@ -274,29 +292,44 @@ int main() { // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]' + +// GEN-AS-NEXT: ArrayInitLoopExpr {{.*}} 'StructWithPointers[2]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'StructWithPointers[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'StructWithPointers[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' '__wrapper_class' +// GEN-AS-NEXT: CXXConstructExpr {{.*}} 'StructWithPointers' 'void (const StructWithPointers &) noexcept' +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'StructWithPointers' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'StructWithPointers *' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'StructWithPointers[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'StructWithPointers[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' '__wrapper_class' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' + +// GLOB-AS-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]' // Initializer for StructWithPointersArray[0] -// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithPointers' 'void (const StructWithPointers &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' -// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: CXXConstructExpr {{.*}} 'StructWithPointers' 'void (const StructWithPointers &) noexcept' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'StructWithPointers' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 + // Initializer for StructWithPointersArray[1] -// CHECK: CXXConstructExpr {{.*}} 'StructWithPointers' 'void (const StructWithPointers &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' -// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS: CXXConstructExpr {{.*}} 'StructWithPointers' 'void (const StructWithPointers &) noexcept' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'StructWithPointers' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 // Check Kernel_Array_Ptrs_2D parameters // CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class) __attribute__((device_kernel))' @@ -310,98 +343,135 @@ int main() { // CHECK-NEXT: InitListExpr // Initializer for ArrayOfPointers_2D -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]' -// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' +// GEN-AS-NEXT: ArrayInitLoopExpr {{.*}} 'int *[2][3]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2][3]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2][3]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GEN-AS-NEXT: ArrayInitLoopExpr {{.*}} 'int *[3]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[3]' lvalue +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'int *[3]' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int *(*)[3]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2][3]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2][3]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'int *' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int **' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[3]' lvalue +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'int *[3]' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int *(*)[3]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2][3]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2][3]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' + +// GLOB-AS-NEXT: InitListExpr {{.*}} 'int *[2][3]' +// GLOB-AS-NEXT: InitListExpr {{.*}} 'int *[3]' + // Initializer for ArrayOfPointers_2D[0][0] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 -// CHECK-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 // Initializer for ArrayOfPointers_2D[0][1] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 -// CHECK-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 // Initializer for ArrayOfPointers_2D[0][2] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 -// CHECK-NEXT: IntegerLiteral {{.*}} 2 - -// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 2 + +// GLOB-AS-NEXT: InitListExpr {{.*}} 'int *[3]' // Initializer for ArrayOfPointers_2D[1][0] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 -// CHECK-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 // Initializer for ArrayOfPointers_2D[1][1] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 -// CHECK-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 // Initializer for ArrayOfPointers_2D[1][2] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 -// CHECK-NEXT: IntegerLiteral {{.*}} 2 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 2 // Initializer for ArrayOfPointers -// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' +// GEN-AS: ArrayInitLoopExpr {{.*}} 'int *[2]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GEN-AS-NEXT: ArraySubscriptExpr {{.*}} 'int *' lvalue +// GEN-AS-NEXT: ImplicitCastExpr {{.*}} 'int **' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GEN-AS-NEXT: ArrayInitIndexExpr {{.*}} '__size_t' + +// GLOB-AS: InitListExpr {{.*}} 'int *[2]' // Initializer for ArrayOfPointers[0] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 // Initializer for ArrayOfPointers[1] -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index 4a9e460fef2b8..06859cc379019 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck --check-prefixes=CHECK,GEN-AS %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump -fsycl-force-global-as-in-kernel-args %s | FileCheck --check-prefixes=CHECK,GLOB-AS %s // This test checks that compiler generates correct initialization for arguments // that have struct or built-in type inside the OpenCL kernel @@ -97,18 +98,21 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int' // Check kernel parameters -// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct) __attribute__((device_kernel))' -// CHECK: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct' +// GLOB-AS: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct) __attribute__((device_kernel))' +// GEN-AS: {{.*}}kernel_struct{{.*}} 'void (test_struct) __attribute__((device_kernel))' +// GLOB-AS: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct' +// GEN-AS: ParmVarDecl {{.*}} used _arg_s 'test_struct' // Check that lambda field of struct type is initialized // CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' // CHECK-NEXT: CXXConstructExpr {{.*}} 'test_struct' 'void (const test_struct &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const test_struct' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'test_struct' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'test_struct *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_test_struct *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'test_struct' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'test_struct *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_test_struct *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_s' // Check kernel parameters // CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class) __attribute__((device_kernel))' @@ -125,30 +129,37 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_data_addr' '__global int *' -// CHECK: InitListExpr {{.*}} 'int *[2]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' -// CHECK-NEXT: IntegerLiteral {{.*}} 0 -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' -// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' -// CHECK-NEXT: IntegerLiteral {{.*}} 1 - -// CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple) __attribute__((device_kernel))' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple' +// GLOB-AS: InitListExpr {{.*}} 'int *[2]' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 0 +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} 'int *' +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// GLOB-AS-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue +// GLOB-AS-NEXT: ImplicitCastExpr {{.*}} '__global int **' +// GLOB-AS-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue . +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' +// GLOB-AS-NEXT: IntegerLiteral {{.*}} 1 +// GEN-AS-NEXT: ArrayInitLoopExpr {{.*}} 'int *[2]' +// GEN-AS-NEXT: OpaqueValueExpr {{.*}} 'int *[2]' lvalue +// GEN-AS-NEXT: MemberExpr {{.*}} 'int *[2]' lvalue . +// GEN-AS-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array' + +// GLOB-AS: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple) __attribute__((device_kernel))' +// GEN-AS: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (Nested::TDS) __attribute__((device_kernel))' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used _arg_tds 'Nested::TDS':'test_struct_simple' // CHECK: VarDecl {{.*}} used __SYCLKernel // CHECK: InitListExpr // CHECK: CXXConstructExpr {{.*}} 'Nested::TDS':'test_struct_simple' 'void (const test_struct_simple &) noexcept' // CHECK: ImplicitCastExpr {{.*}} 'const test_struct_simple' lvalue -// CHECK: UnaryOperator {{.*}} 'Nested::TDS':'test_struct_simple' lvalue prefix '*' cannot overflow -// CHECK: CXXReinterpretCastExpr {{.*}} 'Nested::TDS *' reinterpret_cast -// CHECK: UnaryOperator {{.*}} '__generated_test_struct_simple *' prefix '&' cannot overflow -// CHECK: DeclRefExpr {{.*}} '__generated_test_struct_simple' lvalue ParmVar {{.*}} '_arg_tds' '__generated_test_struct_simple' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Nested::TDS':'test_struct_simple' lvalue ParmVar {{.*}} '_arg_tds' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Nested::TDS':'test_struct_simple' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Nested::TDS *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_test_struct_simple *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_test_struct_simple' lvalue ParmVar {{.*}} '_arg_tds' '__generated_test_struct_simple' diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 91366ebfcde44..edb83284c1d21 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -164,25 +164,25 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return SimpleStructWithPtr.i; }); }); - // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (StructWithPtr) __attribute__((device_kernel))' Nested::TDStrWithPTR TDStructWithPtr; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return TDStructWithPtr.i; }); }); - // CHECK: FunctionDecl {{.*}}TDStr{{.*}} 'void (__generated_StructWithPtr) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}TDStr{{.*}} 'void (Nested::TDStrWithPTR) __attribute__((device_kernel))' StructWithArray t1; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t1.i; }); }); - // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithArray) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (StructWithArray) __attribute__((device_kernel))' DerivedStruct t2; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t2.i; }); }); - // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (__generated_DerivedStruct) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (DerivedStruct) __attribute__((device_kernel))' } { @@ -190,7 +190,7 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return NonTrivialStructWithPtr.i;}); }); - // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (__generated_NonTrivialType) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (NonTrivialType) __attribute__((device_kernel))' NonTrivialType NonTrivialTypeArray[2]{0,0}; myQueue.submit([&](sycl::handler &h) { @@ -202,6 +202,6 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return NonTrivialDerivedStructWithPtr.i;}); }); - // CHECK: FunctionDecl {{.*}}NonTrivialStructInBase{{.*}} 'void (__generated_NonTrivialDerived) __attribute__((device_kernel))' + // CHECK: FunctionDecl {{.*}}NonTrivialStructInBase{{.*}} 'void (NonTrivialDerived) __attribute__((device_kernel))' } } diff --git a/clang/test/SemaSYCL/free_function_array_kernel_param.cpp b/clang/test/SemaSYCL/free_function_array_kernel_param.cpp index 322e495dcabcb..1ce34cae0ab96 100755 --- a/clang/test/SemaSYCL/free_function_array_kernel_param.cpp +++ b/clang/test/SemaSYCL/free_function_array_kernel_param.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \ -// RUN: %s -o - | FileCheck %s +// RUN: %s -o - | FileCheck --check-prefixes=GEN-AS,CHECK %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \ +// RUN: -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck --check-prefixes=GLOB-AS,CHECK %s // This test checks parameter rewriting for free functions with parameters // of type struct with array and array of pointers. @@ -26,15 +28,18 @@ void ff_6(KArgWithPtrArray KArg) { template void ff_6(KArgWithPtrArray KArg); -// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray) __attribute__((device_kernel))' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_KArg '__generated_KArgWithPtrArray' +// GLOB-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray) __attribute__((device_kernel))' +// GEN-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (KArgWithPtrArray<3>) __attribute__((device_kernel))' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_KArg '__generated_KArgWithPtrArray' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_KArg 'KArgWithPtrArray<3>' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: CallExpr // CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(KArgWithPtrArray<3>)' // CHECK-NEXT: DeclRefExpr {{.*}} 'void (KArgWithPtrArray<3>)' lvalue Function {{.*}} 'ff_6' 'void (KArgWithPtrArray<3>)' // CHECK-NEXT: CXXConstructExpr {{.*}} 'KArgWithPtrArray<3>' 'void (const KArgWithPtrArray<3> &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const KArgWithPtrArray<3>' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'KArgWithPtrArray<3>' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'KArgWithPtrArray<3> *' reinterpret_cast *> -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_KArgWithPtrArray *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_KArgWithPtrArray' lvalue ParmVar {{.*}} '__arg_KArg' '__generated_KArgWithPtrArray' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'KArgWithPtrArray<3>' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'KArgWithPtrArray<3> *' reinterpret_cast *> +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_KArgWithPtrArray *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_KArgWithPtrArray' lvalue ParmVar {{.*}} '__arg_KArg' '__generated_KArgWithPtrArray' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'KArgWithPtrArray<3>' lvalue ParmVar {{.*}} '__arg_KArg' diff --git a/clang/test/SemaSYCL/free_function_kernel_params.cpp b/clang/test/SemaSYCL/free_function_kernel_params.cpp index 3f6e27acd8843..5fec7299e4be2 100644 --- a/clang/test/SemaSYCL/free_function_kernel_params.cpp +++ b/clang/test/SemaSYCL/free_function_kernel_params.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \ -// RUN: %s -o - | FileCheck %s +// RUN: %s -o - | FileCheck --check-prefixes=GEN-AS,CHECK %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \ +// RUN: -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck --check-prefixes=GLOB-AS,CHECK %s // This test checks parameter rewriting for free functions with parameters // of type scalar, pointer, non-decomposed struct, work group memory, dynamic work group memory // and special types. @@ -90,10 +92,13 @@ __attribute__((sycl_device)) [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]] void ff_4(NoPointers S1, Pointers S2, Agg S3) { } -// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, __generated_Pointers, __generated_Agg) __attribute__((device_kernel))' +// GEN-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, Pointers, Agg) __attribute__((device_kernel))' +// GLOB-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, __generated_Pointers, __generated_Agg) __attribute__((device_kernel))' // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'NoPointers' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Pointers' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Agg' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 'Pointers' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S3 'Agg' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Pointers' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Agg' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: CallExpr // CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(NoPointers, Pointers, Agg)' @@ -103,25 +108,30 @@ void ff_4(NoPointers S1, Pointers S2, Agg S3) { // CHECK-NEXT: DeclRefExpr {{.*}} 'NoPointers' lvalue ParmVar {{.*}} '__arg_S1' 'NoPointers' // CHECK-NEXT: CXXConstructExpr {{.*}} 'Pointers' 'void (const Pointers &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Pointers' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Pointers' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Pointers *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Pointers *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Pointers' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Pointers' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'Agg' 'void (const Agg &) noexcept' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Pointers' lvalue ParmVar {{.*}} '__arg_S2' 'Pointers' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Pointers' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Pointers *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Pointers *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Pointers' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Pointers' +// CHECK: CXXConstructExpr {{.*}} 'Agg' 'void (const Agg &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Agg' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Agg' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Agg' lvalue ParmVar {{.*}} '__arg_S3' 'Agg' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Agg' __attribute__((sycl_device)) [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]] void ff_5(Agg1 S1, Derived S2, Derived1 S3) { } -// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, __generated_Derived, __generated_Derived1) __attribute__((device_kernel))' +// GEN-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, Derived, Derived1) __attribute__((device_kernel))' +// GLOB-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, __generated_Derived, __generated_Derived1) __attribute__((device_kernel))' // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'Agg1' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Derived1' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 'Derived' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S3 'Derived1' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Derived1' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: CallExpr // CHECK-NEXT: ImplicitCastExpr{{.*}}'void (*)(Agg1, Derived, Derived1)' @@ -131,16 +141,18 @@ void ff_5(Agg1 S1, Derived S2, Derived1 S3) { // CHECK-NEXT: DeclRefExpr {{.*}} 'Agg1' lvalue ParmVar {{.*}} '__arg_S1' 'Agg1' // CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived' 'void (const Derived &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Derived' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Derived' lvalue ParmVar {{.*}} '__arg_S2' 'Derived' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Derived' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Derived *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Derived' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived' +// CHECK: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived1' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Derived1' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Derived1' lvalue ParmVar {{.*}} '__arg_S3' 'Derived1' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S3' '__generated_Derived1' template __attribute__((sycl_device)) @@ -150,9 +162,12 @@ __attribute__((sycl_device)) // Explicit instantiation. template void ff_6(Agg S1, Derived1 S2, int); -// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_Agg, __generated_Derived1, int) __attribute__((device_kernel))' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 '__generated_Agg' -// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived1' +// GEN-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg, Derived1, int) __attribute__((device_kernel))' +// GLOB-AS: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_Agg, __generated_Derived1, int) __attribute__((device_kernel))' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S1 'Agg' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 'Derived1' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S1 '__generated_Agg' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived1' // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_end 'int' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: CallExpr {{.*}} 'void' @@ -160,17 +175,19 @@ template void ff_6(Agg S1, Derived1 S2, int); // CHECK-NEXT: DeclRefExpr {{.*}} 'void (Agg, Derived1, int)' lvalue Function {{.*}} 'ff_6' 'void (Agg, Derived1, int)' // CHECK-NEXT: CXXConstructExpr {{.*}} 'Agg' 'void (const Agg &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Agg' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S1' '__generated_Agg' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Agg' lvalue ParmVar {{.*}} '__arg_S1' 'Agg' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Agg' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Agg *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Agg *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Agg' lvalue ParmVar {{.*}} '__arg_S1' '__generated_Agg' +// CHECK: CXXConstructExpr {{.*}} 'Derived1' 'void (const Derived1 &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const Derived1' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived1' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'Derived1' lvalue ParmVar {{.*}} '__arg_S2' 'Derived1' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'Derived1' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'Derived1 *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_Derived1 *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_Derived1' lvalue ParmVar {{.*}} '__arg_S2' '__generated_Derived1' +// CHECK: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '__arg_end' 'int' __attribute__((sycl_device)) diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index c78c0e40fb4ec..926e25b106996 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s --check-prefixes=CHECK,GEN-AS +// RUN: %clang_cc1 -fsycl-is-device -ast-dump -fsycl-force-global-as-in-kernel-args %s | FileCheck %s --check-prefixes=CHECK,GLOB-AS #include "Inputs/sycl.hpp" @@ -47,12 +48,14 @@ int main() { } // Check declaration of the kernel -// CHECK: derived{{.*}} 'void (base, __generated_second_base, __wrapper_class, +// GEN-AS: derived{{.*}} 'void (base, second_base, __wrapper_class, +// GLOB-AS: derived{{.*}} 'void (base, __generated_second_base, __wrapper_class, // CHECK-SAME: __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg__base 'base' -// CHECK: ParmVarDecl {{.*}} used _arg__base '__generated_second_base' +// GEN-AS: ParmVarDecl {{.*}} used _arg__base 'second_base' +// GLOB-AS: ParmVarDecl {{.*}} used _arg__base '__generated_second_base' // CHECK: ParmVarDecl {{.*}} used _arg_d '__wrapper_class' // CHECK: ParmVarDecl {{.*}} used _arg_AccField '__global char *' // CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' @@ -72,15 +75,15 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg__base' 'base' -// second_base contains pointers and therefore the ParamVar is a new generated -// type. Perform a copy of the corresponding kernel parameter via -// reinterpret_cast. +// second_base contains pointers. With GLOB-AS, the ParamVar is a new generated +// type and a reinterpret_cast is used. With GEN-AS, it is passed directly. // CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base' 'void (const second_base &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const second_base' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'second_base' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'second_base *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_second_base *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '__generated_second_base' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'second_base' lvalue ParmVar {{.*}} '_arg__base' 'second_base' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'second_base' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'second_base *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_second_base *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '__generated_second_base' // third_base contains special type accessor. Therefore it is decomposed and it's // data members are copied from corresponding ParamVar diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index 92f19374818ed..8ccce15c5df51 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -1,10 +1,13 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device \ // RUN: -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml -// RUN: FileCheck -check-prefix=SPIR --input-file %t-host.yaml %s +// RUN: FileCheck --check-prefixes=SPIR,GEN-AS --input-file %t-host.yaml %s +// RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device \ +// RUN: -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml -fsycl-force-global-as-in-kernel-args +// RUN: FileCheck --check-prefixes=SPIR,GLOB-AS --input-file %t-host.yaml %s // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fsycl-is-device \ // RUN: -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml -// RUN: FileCheck -check-prefix=NVPTX --input-file %t-host.yaml %s +// RUN: FileCheck --check-prefixes=NVPTX --input-file %t-host.yaml %s // The test generates remarks about the kernel argument, their location and type // in the resulting yaml file. @@ -69,7 +72,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -90,7 +93,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -111,7 +114,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -132,7 +135,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -153,7 +156,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -174,7 +177,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -195,7 +198,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -216,7 +219,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -237,7 +240,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -258,7 +261,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -279,7 +282,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -300,7 +303,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -321,19 +324,21 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' // SPIR-NEXT: Argument: '12' // SPIR-NEXT: String: ':' -// SPIR-NEXT: String: Compiler generated argument for base class with pointer, -// SPIR-NEXT: String: StructWithPointer +// GEN-AS: String: Compiler generated argument for base class, +// GLOB-AS: String: Compiler generated argument for base class with pointer, +// SPIR: String: StructWithPointer // SPIR-NEXT: String: ' (' // SPIR-NEXT: String: '' // SPIR-NEXT: String: 'Type:' -// SPIR-NEXT: String: Compiler generated -// SPIR-NEXT: String: ', ' +// GEN-AS: String: StructWithPointer +// GLOB-AS: String: Compiler generated +// SPIR: String: ', ' // SPIR-NEXT: String: 'Size: ' // SPIR-NEXT: Argument: '8' // SPIR-NEXT: String: ')' @@ -342,7 +347,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -363,7 +368,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -384,7 +389,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -405,7 +410,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -426,19 +431,22 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Line: 36, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' // SPIR-NEXT: Argument: '17' // SPIR-NEXT: String: ':' -// SPIR-NEXT: String: Compiler generated argument for object with pointer, -// SPIR-NEXT: String: Obj -// SPIR-NEXT: String: ' (' -// SPIR-NEXT: String: '' -// SPIR-NEXT: String: 'Type:' -// SPIR-NEXT: String: Compiler generated -// SPIR-NEXT: String: ', ' +// GEN-AS: String: 'Field:Obj, ' +// GEN-AS-NEXT: String: 'Type:' +// GEN-AS-NEXT: String: StructWithPointer +// GLOB-AS: String: Compiler generated argument for object with pointer, +// GLOB-AS-NEXT: String: Obj +// GLOB-AS-NEXT: String: ' (' +// GLOB-AS-NEXT: String: '' +// GLOB-AS-NEXT: String: 'Type:' +// GLOB-AS-NEXT: String: Compiler generated +// SPIR: String: ', ' // SPIR-NEXT: String: 'Size: ' // SPIR-NEXT: Argument: '8' // SPIR-NEXT: String: ')' @@ -448,7 +456,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 59, Column: 9 } +// SPIR-NEXT: Line: 62, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -469,7 +477,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 59, Column: 9 } +// SPIR-NEXT: Line: 62, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -490,7 +498,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 59, Column: 9 } +// SPIR-NEXT: Line: 62, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -511,7 +519,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 59, Column: 9 } +// SPIR-NEXT: Line: 62, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -532,7 +540,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 59, Column: 9 } +// SPIR-NEXT: Line: 62, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -553,7 +561,7 @@ int main() { // NVPTX: Pass:{{.*}}sycl // NVPTX: Name:{{.*}}Region // NVPTX: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// NVPTX: Line: 59, Column: 9 } +// NVPTX: Line: 62, Column: 9 } // NVPTX-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // NVPTX-NEXT: Args: // NVPTX-NEXT: String: 'Arg ' diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 3702fe2293304..e820e929e7203 100644 --- a/clang/test/SemaSYCL/union-kernel-param2.cpp +++ b/clang/test/SemaSYCL/union-kernel-param2.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck --check-prefixes=CHECK,GEN-AS %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -fsycl-force-global-as-in-kernel-args %s | FileCheck --check-prefixes=CHECK,GLOB-AS %s // This test checks that compiler generates correct kernel arguments for // a struct-with-an-array-of-unions and a array-of-struct-with-a-union. @@ -95,11 +96,13 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField // CHECK-NEXT: MemberExpr {{.*}} lvalue .struct_mem -// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp{{.*}})' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp{{.*}})' // Check kernel_C parameters -// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr) __attribute__((device_kernel))' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '__generated_MyStructWithPtr' +// GEN-AS: FunctionDecl {{.*}}kernel_C{{.*}} 'void (struct MyStructWithPtr) __attribute__((device_kernel))' +// GEN-AS-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem 'struct MyStructWithPtr' +// GLOB-AS: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr) __attribute__((device_kernel))' +// GLOB-AS-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '__generated_MyStructWithPtr' // Check kernel_C inits // CHECK-NEXT: CompoundStmt @@ -108,7 +111,8 @@ int main() { // CHECK-NEXT: InitListExpr // CHECK-NEXT: CXXConstructExpr {{.*}} 'struct MyStructWithPtr' 'void (const MyStructWithPtr &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyStructWithPtr' lvalue -// CHECK-NEXT: UnaryOperator {{.*}} 'struct MyStructWithPtr' lvalue prefix '*' cannot overflow -// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'struct MyStructWithPtr *' reinterpret_cast -// CHECK-NEXT: UnaryOperator {{.*}} '__generated_MyStructWithPtr *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '__generated_MyStructWithPtr' +// GEN-AS-NEXT: DeclRefExpr {{.*}} 'struct MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' 'struct MyStructWithPtr' +// GLOB-AS-NEXT: UnaryOperator {{.*}} 'struct MyStructWithPtr' lvalue prefix '*' cannot overflow +// GLOB-AS-NEXT: CXXReinterpretCastExpr {{.*}} 'struct MyStructWithPtr *' reinterpret_cast +// GLOB-AS-NEXT: UnaryOperator {{.*}} '__generated_MyStructWithPtr *' prefix '&' cannot overflow +// GLOB-AS-NEXT: DeclRefExpr {{.*}} '__generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '__generated_MyStructWithPtr'