diff --git a/clang/docs/ClangCommandLineReference.rst b/clang/docs/ClangCommandLineReference.rst --- a/clang/docs/ClangCommandLineReference.rst +++ b/clang/docs/ClangCommandLineReference.rst @@ -2147,6 +2147,13 @@ OpenCL only. Allow unsafe floating-point optimizations. Also implies -cl-no-signed-zeros and -cl-mad-enable. +C++ for OpenCL flags +-------------------- + +.. option:: -clcxx-clc-compatibility + +C++ for OpenCL only. Accept OpenCL C constructs that would otherwise be undefined in C++ for OpenCL. Favours compatibility with OpenCL C rather than with C++. + Target-dependent compilation options ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ .. option:: -G, -G=, -msmall-data-threshold= diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -202,6 +202,7 @@ LANGOPT(OpenCL , 1, 0, "OpenCL") LANGOPT(OpenCLVersion , 32, 0, "OpenCL C version") LANGOPT(OpenCLCPlusPlus , 1, 0, "C++ for OpenCL") +LANGOPT(OpenCLCPlusPlusCLCCompat , 1, 0, "C++ for OpenCL - OpenCL C compatibility mode") LANGOPT(OpenCLCPlusPlusVersion , 32, 0, "C++ for OpenCL version") LANGOPT(NativeHalfType , 1, 0, "Native half type support") LANGOPT(NativeHalfArgsAndReturns, 1, 0, "Native half args and returns") diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def --- a/clang/include/clang/Basic/TokenKinds.def +++ b/clang/include/clang/Basic/TokenKinds.def @@ -287,7 +287,7 @@ KEYWORD(int , KEYALL) KEYWORD(long , KEYALL) KEYWORD(register , KEYALL) -KEYWORD(restrict , KEYC99) +KEYWORD(restrict , KEYC99|KEYOPENCLCXXCOMPAT) KEYWORD(return , KEYALL) KEYWORD(short , KEYALL) KEYWORD(signed , KEYALL) diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -124,6 +124,9 @@ def opencl_Group : OptionGroup<"">, Group, DocName<"OpenCL flags">; +def openclcxx_Group : OptionGroup<"">, Group, + DocName<"C++ for OpenCL flags">; + def m_Group : OptionGroup<"">, Group, DocName<"Target-dependent compilation options">; @@ -529,6 +532,8 @@ HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">; def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group, Flags<[CC1Option]>, HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">; +def clcxx_clc_compatibility : Flag<["-"], "clcxx-clc-compatibility">, Group, Flags<[CC1Option]>, + HelpText<"C++ for OpenCL only. Accept OpenCL C constructs that would otherwise be undefined in C++ for OpenCL. Favours compatibility with OpenCL C rather than with C++.">; def client__name : JoinedOrSeparate<["-"], "client_name">; def combine : Flag<["-", "--"], "combine">, Flags<[DriverOption, Unsupported]>; def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">; diff --git a/clang/lib/Basic/IdentifierTable.cpp b/clang/lib/Basic/IdentifierTable.cpp --- a/clang/lib/Basic/IdentifierTable.cpp +++ b/clang/lib/Basic/IdentifierTable.cpp @@ -76,30 +76,31 @@ namespace { enum { - KEYC99 = 0x1, - KEYCXX = 0x2, - KEYCXX11 = 0x4, - KEYGNU = 0x8, - KEYMS = 0x10, - BOOLSUPPORT = 0x20, - KEYALTIVEC = 0x40, - KEYNOCXX = 0x80, - KEYBORLAND = 0x100, - KEYOPENCLC = 0x200, - KEYC11 = 0x400, - KEYNOMS18 = 0x800, - KEYNOOPENCL = 0x1000, - WCHARSUPPORT = 0x2000, - HALFSUPPORT = 0x4000, - CHAR8SUPPORT = 0x8000, - KEYCONCEPTS = 0x10000, - KEYOBJC = 0x20000, - KEYZVECTOR = 0x40000, - KEYCOROUTINES = 0x80000, - KEYMODULES = 0x100000, - KEYCXX2A = 0x200000, - KEYOPENCLCXX = 0x400000, - KEYMSCOMPAT = 0x800000, + KEYC99 = 0x1, + KEYCXX = 0x2, + KEYCXX11 = 0x4, + KEYGNU = 0x8, + KEYMS = 0x10, + BOOLSUPPORT = 0x20, + KEYALTIVEC = 0x40, + KEYNOCXX = 0x80, + KEYBORLAND = 0x100, + KEYOPENCLC = 0x200, + KEYC11 = 0x400, + KEYNOMS18 = 0x800, + KEYNOOPENCL = 0x1000, + WCHARSUPPORT = 0x2000, + HALFSUPPORT = 0x4000, + CHAR8SUPPORT = 0x8000, + KEYCONCEPTS = 0x10000, + KEYOBJC = 0x20000, + KEYZVECTOR = 0x40000, + KEYCOROUTINES = 0x80000, + KEYMODULES = 0x100000, + KEYCXX2A = 0x200000, + KEYOPENCLCXX = 0x400000, + KEYMSCOMPAT = 0x800000, + KEYOPENCLCXXCOMPAT = 0x1000000, KEYALLCXX = KEYCXX | KEYCXX11 | KEYCXX2A, KEYALL = (0xffffff & ~KEYNOMS18 & ~KEYNOOPENCL) // KEYNOMS18 and KEYNOOPENCL are used to exclude. @@ -137,6 +138,8 @@ if (LangOpts.OpenCL && !LangOpts.OpenCLCPlusPlus && (Flags & KEYOPENCLC)) return KS_Enabled; if (LangOpts.OpenCLCPlusPlus && (Flags & KEYOPENCLCXX)) return KS_Enabled; + if (LangOpts.OpenCLCPlusPlusCLCCompat && (Flags & KEYOPENCLCXXCOMPAT)) + return KS_Extension; if (!LangOpts.CPlusPlus && (Flags & KEYNOCXX)) return KS_Enabled; if (LangOpts.C11 && (Flags & KEYC11)) return KS_Enabled; // We treat bridge casts as objective-C keywords so we can warn on them diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -2620,7 +2620,8 @@ options::OPT_cl_no_signed_zeros, options::OPT_cl_denorms_are_zero, options::OPT_cl_fp32_correctly_rounded_divide_sqrt, - options::OPT_cl_uniform_work_group_size + options::OPT_cl_uniform_work_group_size, + options::OPT_clcxx_clc_compatibility }; if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) { diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2522,6 +2522,10 @@ << Args.getLastArg(OPT_cl_strict_aliasing)->getAsString(Args); } + if (Args.getLastArg(OPT_clcxx_clc_compatibility)) { + Opts.OpenCLCPlusPlusCLCCompat = true; + } + // We abuse '-f[no-]gnu-keywords' to force overriding all GNU-extension // keywords. This behavior is provided by GCC's poorly named '-fasm' flag, // while a subset (the non-C++ GNU keywords) is provided by GCC's diff --git a/clang/test/CodeGenOpenCLCXX/compat-clc-kernel-arg-info.cl b/clang/test/CodeGenOpenCLCXX/compat-clc-kernel-arg-info.cl new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenOpenCLCXX/compat-clc-kernel-arg-info.cl @@ -0,0 +1,130 @@ +// RUN: %clang_cc1 %s -cl-std=clc++ -clcxx-clc-compatibility -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=clc++ -clcxx-clc-compatibility -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO + +kernel void foo(global int * globalintp, global int * restrict globalintrestrictp, + global const int * globalconstintp, + global const int * restrict globalconstintrestrictp, + constant int * constantintp, constant int * restrict constantintrestrictp, + global const volatile int * globalconstvolatileintp, + global const volatile int * restrict globalconstvolatileintrestrictp, + global volatile int * globalvolatileintp, + global volatile int * restrict globalvolatileintrestrictp, + local int * localintp, local int * restrict localintrestrictp, + local const int * localconstintp, + local const int * restrict localconstintrestrictp, + local const volatile int * localconstvolatileintp, + local const volatile int * restrict localconstvolatileintrestrictp, + local volatile int * localvolatileintp, + local volatile int * restrict localvolatileintrestrictp, + int X, const int constint, const volatile int constvolatileint, + volatile int volatileint) { + *globalintrestrictp = constint + volatileint; +} +// CHECK: define spir_kernel void @foo{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD11:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD12:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD13:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD13]] +// CHECK: !kernel_arg_type_qual ![[MD14:[0-9]+]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]] + +kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3, read_write image1d_t img4) { +} +// CHECK: define spir_kernel void @foo2{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD22:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD23:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD23]] +// CHECK: !kernel_arg_type_qual ![[MD24:[0-9]+]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD25:[0-9]+]] + +kernel void foo3(__global half * X) { +} +// CHECK: define spir_kernel void @foo3{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD31:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD32:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD33:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD33]] +// CHECK: !kernel_arg_type_qual ![[MD34:[0-9]+]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD35:[0-9]+]] + +typedef unsigned int myunsignedint; +kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) { +} +// CHECK: define spir_kernel void @foo4{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD42:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD43:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD44:[0-9]+]] +// CHECK: !kernel_arg_type_qual ![[MD45:[0-9]+]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD46:[0-9]+]] + +typedef image1d_t myImage; +kernel void foo5(myImage img1, write_only image1d_t img2) { +} +// CHECK: define spir_kernel void @foo5{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD51:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD52:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD53:[0-9]+]] +// CHECK: !kernel_arg_type_qual ![[MD45]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]] + +typedef char char16 __attribute__((ext_vector_type(16))); +__kernel void foo6(__global char16 arg[]) {} +// CHECK: !kernel_arg_type ![[MD61:[0-9]+]] +// ARGINFO: !kernel_arg_name ![[MD62:[0-9]+]] + +typedef read_only image1d_t ROImage; +typedef write_only image1d_t WOImage; +typedef read_write image1d_t RWImage; +kernel void foo7(ROImage ro, WOImage wo, RWImage rw) { +} +// CHECK: define spir_kernel void @foo7{{[^!]+}} +// CHECK: !kernel_arg_addr_space ![[MD71:[0-9]+]] +// CHECK: !kernel_arg_access_qual ![[MD72:[0-9]+]] +// CHECK: !kernel_arg_type ![[MD73:[0-9]+]] +// CHECK: !kernel_arg_base_type ![[MD74:[0-9]+]] +// CHECK: !kernel_arg_type_qual ![[MD75:[0-9]+]] +// CHECK-NOT: !kernel_arg_name +// ARGINFO: !kernel_arg_name ![[MD76:[0-9]+]] + +// CHECK: ![[MD11]] = !{i32 1, i32 1, i32 1, i32 1, i32 2, i32 2, i32 1, i32 1, i32 1, i32 1, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 0, i32 0, i32 0, i32 0} +// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"} +// CHECK: ![[MD13]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int", !"int", !"int", !"int"} +// CHECK: ![[MD14]] = !{!"", !"restrict", !"const", !"restrict const", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"", !"", !""} +// ARGINFO: ![[MD15]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"constantintp", !"constantintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp", !"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp", !"X", !"constint", !"constvolatileint", !"volatileint"} +// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1, i32 1} +// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only", !"read_write"} +// CHECK: ![[MD23]] = !{!"image1d_t", !"image2d_t", !"image2d_array_t", !"image1d_t"} +// CHECK: ![[MD24]] = !{!"", !"", !"", !""} +// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3", !"img4"} +// CHECK: ![[MD31]] = !{i32 1} +// CHECK: ![[MD32]] = !{!"none"} +// CHECK: ![[MD33]] = !{!"half*"} +// CHECK: ![[MD34]] = !{!""} +// ARGINFO: ![[MD35]] = !{!"X"} +// CHECK: ![[MD41]] = !{i32 1, i32 1} +// CHECK: ![[MD42]] = !{!"none", !"none"} +// CHECK: ![[MD43]] = !{!"uint*", !"myunsignedint*"} +// CHECK: ![[MD44]] = !{!"uint*", !"uint*"} +// CHECK: ![[MD45]] = !{!"", !""} +// ARGINFO: ![[MD46]] = !{!"X", !"Y"} +// CHECK: ![[MD51]] = !{!"read_only", !"write_only"} +// CHECK: ![[MD52]] = !{!"myImage", !"image1d_t"} +// CHECK: ![[MD53]] = !{!"image1d_t", !"image1d_t"} +// ARGINFO: ![[MD54]] = !{!"img1", !"img2"} +// CHECK: ![[MD61]] = !{!"char16*"} +// ARGINFO: ![[MD62]] = !{!"arg"} +// CHECK: ![[MD71]] = !{i32 1, i32 1, i32 1} +// CHECK: ![[MD72]] = !{!"read_only", !"write_only", !"read_write"} +// CHECK: ![[MD73]] = !{!"ROImage", !"WOImage", !"RWImage"} +// CHECK: ![[MD74]] = !{!"image1d_t", !"image1d_t", !"image1d_t"} +// CHECK: ![[MD75]] = !{!"", !"", !""} +// ARGINFO: ![[MD76]] = !{!"ro", !"wo", !"rw"} +