diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -1784,7 +1784,7 @@ case LangAS::opencl_local: return "__local"; case LangAS::opencl_private: - return ""; + return "__private"; case LangAS::opencl_constant: return "__constant"; case LangAS::opencl_generic: diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -8381,7 +8381,7 @@ QualType DesugaredTy = Ty; do { ArrayRef Names(SizeTypeNames); - auto Match = llvm::find(Names, DesugaredTy.getAsString()); + auto Match = llvm::find(Names, DesugaredTy.getUnqualifiedType().getAsString()); if (Names.end() != Match) return true; diff --git a/clang/test/AST/language_address_space_attribute.cpp b/clang/test/AST/language_address_space_attribute.cpp --- a/clang/test/AST/language_address_space_attribute.cpp +++ b/clang/test/AST/language_address_space_attribute.cpp @@ -22,10 +22,10 @@ // CHECK: VarDecl {{.*}} z_constant '__constant int *' [[clang::opencl_constant]] int *z_constant; - // CHECK: VarDecl {{.*}} x_private 'int *' + // CHECK: VarDecl {{.*}} x_private '__private int *' __attribute__((opencl_private)) int *x_private; - // CHECK: VarDecl {{.*}} z_private 'int *' + // CHECK: VarDecl {{.*}} z_private '__private int *' [[clang::opencl_private]] int *z_private; // CHECK: VarDecl {{.*}} x_generic '__generic int *' diff --git a/clang/test/Index/opencl-types.cl b/clang/test/Index/opencl-types.cl --- a/clang/test/Index/opencl-types.cl +++ b/clang/test/Index/opencl-types.cl @@ -16,12 +16,12 @@ double4 vectorDouble; } -// CHECK: VarDecl=scalarHalf:11:8 (Definition){{( \(invalid\))?}} [type=half] [typekind=Half] [isPOD=1] -// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=half4] [typekind=Typedef] [canonicaltype=half __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] -// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=float] [typekind=Float] [isPOD=1] -// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=float4] [typekind=Typedef] [canonicaltype=float __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] -// CHECK: VarDecl=scalarDouble:15:10 (Definition){{( \(invalid\))?}} [type=double] [typekind=Double] [isPOD=1] -// CHECK: VarDecl=vectorDouble:16:11 (Definition){{( \(invalid\))?}} [type=double4] [typekind=Typedef] [canonicaltype=double __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] +// CHECK: VarDecl=scalarHalf:11:8 (Definition){{( \(invalid\))?}} [type=__private half] [typekind=Half] [isPOD=1] +// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=__private half4] [typekind=Typedef] [canonicaltype=half __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] +// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=__private float] [typekind=Float] [isPOD=1] +// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=__private float4] [typekind=Typedef] [canonicaltype=float __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] +// CHECK: VarDecl=scalarDouble:15:10 (Definition){{( \(invalid\))?}} [type=__private double] [typekind=Double] [isPOD=1] +// CHECK: VarDecl=vectorDouble:16:11 (Definition){{( \(invalid\))?}} [type=__private double4] [typekind=Typedef] [canonicaltype=double __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1] #pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable @@ -38,18 +38,18 @@ void kernel OCLImage2dArrayMSAADepthROTest(read_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRO); void kernel OCLImage3dROTest(read_only image3d_t scalarOCLImage3dRO); -// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition){{( \(invalid\))?}} [type=__read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition){{( \(invalid\))?}} [type=__read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition){{( \(invalid\))?}} [type=__read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition){{( \(invalid\))?}} [type=__read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__private __read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__private __read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__private __read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__private __read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__private __read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__private __read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__private __read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__private __read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1] void kernel OCLImage1dWOTest(write_only image1d_t scalarOCLImage1dWO); void kernel OCLImage1dArrayWOTest(write_only image1d_array_t scalarOCLImage1dArrayWO); @@ -64,18 +64,18 @@ void kernel OCLImage2dArrayMSAADepthWOTest(write_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthWO); void kernel OCLImage3dWOTest(write_only image3d_t scalarOCLImage3dWO); -// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition){{( \(invalid\))?}} [type=__write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition){{( \(invalid\))?}} [type=__write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition){{( \(invalid\))?}} [type=__write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition){{( \(invalid\))?}} [type=__write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition){{( \(invalid\))?}} [type=__write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__private __write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__private __write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__private __write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__private __write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__private __write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__private __write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__private __write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition){{( \(invalid\))?}} [type=__private __write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1] void kernel OCLImage1dRWTest(read_write image1d_t scalarOCLImage1dRW); void kernel OCLImage1dArrayRWTest(read_write image1d_array_t scalarOCLImage1dArrayRW); @@ -90,24 +90,24 @@ void kernel OCLImage2dArrayMSAADepthRWTest(read_write image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRW); void kernel OCLImage3dRWTest(read_write image3d_t scalarOCLImage3dRW); -// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition){{( \(invalid\))?}} [type=__read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition){{( \(invalid\))?}} [type=__read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition){{( \(invalid\))?}} [type=__read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition){{( \(invalid\))?}} [type=__read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1] -// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__private __read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__private __read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__private __read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__private __read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__private __read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__private __read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__private __read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1] +// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__private __read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1] void kernel intPipeTestRO(read_only pipe int scalarPipe); void kernel intPipeTestWO(write_only pipe int scalarPipe); -// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=read_only pipe int] [typekind=Pipe] [isPOD=0] -// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=write_only pipe int] [typekind=Pipe] [isPOD=0] +// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=__private read_only pipe int] [typekind=Pipe] [isPOD=0] +// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=__private write_only pipe int] [typekind=Pipe] [isPOD=0] #define CLK_ADDRESS_CLAMP_TO_EDGE 2 #define CLK_NORMALIZED_COORDS_TRUE 1 @@ -121,9 +121,9 @@ } // CHECK: VarDecl=scalarOCLSampler:117:19 (Definition) [type=const sampler_t] [typekind=Typedef] const [canonicaltype=const sampler_t] [canonicaltypekind=OCLSampler] [isPOD=1] -// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=clk_event_t] [typekind=Typedef] [canonicaltype=clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1] -// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=queue_t] [typekind=Typedef] [canonicaltype=queue_t] [canonicaltypekind=OCLQueue] [isPOD=1] -// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=reserve_id_t] [typekind=Typedef] [canonicaltype=reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1] +// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=__private clk_event_t] [typekind=Typedef] [canonicaltype=__private clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1] +// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=__private queue_t] [typekind=Typedef] [canonicaltype=__private queue_t] [canonicaltypekind=OCLQueue] [isPOD=1] +// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=__private reserve_id_t] [typekind=Typedef] [canonicaltype=__private reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1] #pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable @@ -131,4 +131,4 @@ intel_sub_group_avc_mce_payload_t mce_payload; } -// CHECK: VarDecl=mce_payload:131:37 (Definition){{( \(invalid\))?}} [type=intel_sub_group_avc_mce_payload_t] [typekind=Typedef] [canonicaltype=intel_sub_group_avc_mce_payload_t] [canonicaltypekind=OCLIntelSubgroupAVCMcePayload] [isPOD=1] +// CHECK: VarDecl=mce_payload:131:37 (Definition){{( \(invalid\))?}} [type=__private intel_sub_group_avc_mce_payload_t] [typekind=Typedef] [canonicaltype=__private intel_sub_group_avc_mce_payload_t] [canonicaltypekind=OCLIntelSubgroupAVCMcePayload] [isPOD=1] diff --git a/clang/test/Parser/opencl-astype.cl b/clang/test/Parser/opencl-astype.cl --- a/clang/test/Parser/opencl-astype.cl +++ b/clang/test/Parser/opencl-astype.cl @@ -11,7 +11,7 @@ typedef __attribute__(( ext_vector_type(4) )) double double4; float4 f4; - double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and 'float4' (vector of 4 'float' values) must match}} + double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and '__private float4' (vector of 4 'float' values) must match}} // Verify int4->float3, float3->int4 works. int4 i4; diff --git a/clang/test/Parser/opencl-atomics-cl20.cl b/clang/test/Parser/opencl-atomics-cl20.cl --- a/clang/test/Parser/opencl-atomics-cl20.cl +++ b/clang/test/Parser/opencl-atomics-cl20.cl @@ -66,9 +66,9 @@ atomic_int i; foo(&i); // OpenCL v2.0 s6.13.11.8, arithemtic operations are not permitted on atomic types. - i++; // expected-error {{invalid argument type 'atomic_int' (aka '_Atomic(int)') to unary expression}} + i++; // expected-error {{invalid argument type '__private atomic_int' (aka '__private _Atomic(int)') to unary expression}} i = 1; // expected-error {{atomic variable can be assigned to a variable only in global address space}} - i += 1; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'int')}} - i = i + i; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'atomic_int')}} + i += 1; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and 'int')}} + i = i + i; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and '__private atomic_int')}} } #endif diff --git a/clang/test/SemaOpenCL/access-qualifier.cl b/clang/test/SemaOpenCL/access-qualifier.cl --- a/clang/test/SemaOpenCL/access-qualifier.cl +++ b/clang/test/SemaOpenCL/access-qualifier.cl @@ -25,11 +25,11 @@ kernel void k1(img1d_wo img) { - myRead(img); // expected-error {{passing 'img1d_wo' (aka '__write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}} + myRead(img); // expected-error {{passing '__private img1d_wo' (aka '__private __write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}} } kernel void k2(img1d_ro img) { - myWrite(img); // expected-error {{passing 'img1d_ro' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}} + myWrite(img); // expected-error {{passing '__private img1d_ro' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}} } kernel void k3(img1d_wo img) { @@ -43,7 +43,7 @@ #endif kernel void k5(img1d_ro_default img) { - myWrite(img); // expected-error {{passing 'img1d_ro_default' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}} + myWrite(img); // expected-error {{passing '__private img1d_ro_default' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}} } kernel void k6(img1d_ro img) { @@ -71,7 +71,7 @@ #if __OPENCL_C_VERSION__ >= 200 void myPipeWrite(write_only pipe int); // expected-note {{passing argument to parameter here}} kernel void k14(read_only pipe int p) { - myPipeWrite(p); // expected-error {{passing 'read_only pipe int' to parameter of incompatible type 'write_only pipe int'}} + myPipeWrite(p); // expected-error {{passing '__private read_only pipe int' to parameter of incompatible type 'write_only pipe int'}} } #endif @@ -93,7 +93,7 @@ // expected-note@-2 {{previously declared 'read_only' here}} kernel void pass_ro_typedef_to_wo(ROPipeInt p) { - myPipeWrite(p); // expected-error {{passing 'ROPipeInt' (aka 'read_only pipe int') to parameter of incompatible type 'write_only pipe int'}} + myPipeWrite(p); // expected-error {{passing '__private ROPipeInt' (aka '__private read_only pipe int') to parameter of incompatible type 'write_only pipe int'}} // expected-note@-25 {{passing argument to parameter here}} } #endif diff --git a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl --- a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl +++ b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl @@ -66,7 +66,7 @@ #if !__OPENCL_CPP_VERSION__ // expected-note@-2{{passing argument to parameter 'arg_priv' here}} #else -// expected-note-re@-4{{candidate function not viable: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to default address space in 1st argument}} +// expected-note-re@-4{{candidate function not viable: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to address space '__private' in 1st argument}} #endif void f_gen(__generic int *arg_gen) {} @@ -85,45 +85,45 @@ AS int *var_init1 = arg_glob; #ifdef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error@-3{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}} +// expected-error@-3{{initializing '__constant int *__private' with an expression of type '__global int *__private' changes address space of pointer}} #else -// expected-error@-5{{cannot initialize a variable of type '__constant int *' with an lvalue of type '__global int *'}} +// expected-error@-5{{cannot initialize a variable of type '__constant int *__private' with an lvalue of type '__global int *__private'}} #endif #endif AS int *var_init2 = arg_loc; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}} +// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__local int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__local int *'}} +// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__local int *__private'}} #endif #endif AS int *var_init3 = arg_const; #ifndef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}} +// expected-error-re@-3{{initializing '__{{global|generic}} int *__private' with an expression of type '__constant int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{cannot initialize a variable of type '__{{global|generic}} int *' with an lvalue of type '__constant int *'}} +// expected-error-re@-5{{cannot initialize a variable of type '__{{global|generic}} int *__private' with an lvalue of type '__constant int *__private'}} #endif #endif AS int *var_init4 = arg_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}} +// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__private int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type 'int *'}} +// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__private int *__private'}} #endif #endif AS int *var_init5 = arg_gen; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}} +// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__generic int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__generic int *'}} +// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__generic int *__private'}} #endif #endif @@ -157,9 +157,9 @@ AS int *var_cast4 = (AS int *)arg_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}} +// expected-error-re@-3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}} #else -// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}} +// expected-error-re@-5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}} #endif #endif @@ -176,45 +176,45 @@ var_impl = arg_glob; #ifdef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error@-3{{assigning '__global int *' to '__constant int *' changes address space of pointer}} +// expected-error@-3{{assigning '__global int *__private' to '__constant int *__private' changes address space of pointer}} #else -// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *'}} +// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *__private'}} #endif #endif var_impl = arg_loc; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}} +// expected-error-re@-3{{assigning '__local int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *'}} +// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *__private'}} #endif #endif var_impl = arg_const; #ifndef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}} +// expected-error-re@-3{{assigning '__constant int *__private' to '__{{global|generic}} int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *'}} +// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *__private'}} #endif #endif var_impl = arg_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}} +// expected-error-re@-3{{assigning '__private int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type 'int *'}} +// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__private int *__private'}} #endif #endif var_impl = arg_gen; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}} +// expected-error-re@-3{{assigning '__generic int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}} #else -// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *'}} +// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *__private'}} #endif #endif @@ -248,9 +248,9 @@ var_cast4 = (AS int *)arg_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}} +// expected-error-re@-3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}} #else -// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}} +// expected-error-re@-5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}} #endif #endif @@ -294,9 +294,9 @@ b = var_cmp <= arg_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{comparison between ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}} +// expected-error-re@-3{{comparison between ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}} #else -// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and 'int *')}} +// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and '__private int *')}} #endif #endif @@ -327,7 +327,7 @@ b = var_sub - arg_priv; #ifndef GENERIC -// expected-error-re@-2{{arithmetic operation with operands of type ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}} +// expected-error-re@-2{{arithmetic operation with operands of type ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}} #endif b = var_sub - arg_gen; @@ -338,7 +338,7 @@ f_glob(var_sub); #ifndef GLOBAL #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}} +// expected-error-re@-3{{passing '__{{constant|generic}} int *__private' to parameter of type '__global int *' changes address space of pointer}} #else // expected-error@-5{{no matching function for call to 'f_glob'}} #endif @@ -346,7 +346,7 @@ f_loc(var_sub); #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}} +// expected-error-re@-2{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__local int *' changes address space of pointer}} #else // expected-error@-4{{no matching function for call to 'f_loc'}} #endif @@ -354,7 +354,7 @@ f_const(var_sub); #ifndef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}} +// expected-error-re@-3{{passing '__{{global|generic}} int *__private' to parameter of type '__constant int *' changes address space of pointer}} #else // expected-error@-5{{no matching function for call to 'f_const'}} #endif @@ -362,7 +362,7 @@ f_priv(var_sub); #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}} +// expected-error-re@-2{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__private int *' changes address space of pointer}} #else // expected-error@-4{{no matching function for call to 'f_priv'}} #endif @@ -370,7 +370,7 @@ f_gen(var_sub); #ifdef CONSTANT #if !__OPENCL_CPP_VERSION__ -// expected-error@-3{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}} +// expected-error@-3{{passing '__constant int *__private' to parameter of type '__generic int *' changes address space of pointer}} #else // expected-error@-5{{no matching function for call to 'f_gen'}} #endif @@ -414,9 +414,9 @@ var_gen = 0 ? var_cond : var_priv; #ifndef GENERIC #if !__OPENCL_CPP_VERSION__ -// expected-error-re@-3{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}} +// expected-error-re@-3{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}} #else -// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and 'int *')}} +// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and '__private int *')}} #endif #endif @@ -470,12 +470,12 @@ __private char *var_priv_ch; var_void_gen = 0 ? var_cond : var_priv_ch; #if __OPENCL_CPP_VERSION__ -// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and 'char *')}} +// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__private char *')}} #else #ifndef GENERIC -// expected-error-re@-5{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}} +// expected-error-re@-5{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and '__private char *') which are pointers to non-overlapping address spaces}} #else -// expected-warning@-7{{pointer type mismatch ('__generic int *' and 'char *')}} +// expected-warning@-7{{pointer type mismatch ('__generic int *' and '__private char *')}} #endif #endif @@ -517,7 +517,7 @@ var_as_as_int = var_asc_asc_int; #if !__OPENCL_CPP_VERSION__ #ifdef GENERIC -// expected-error@-3 {{assigning '__local int *__local *' to '__generic int *__generic *' changes address space of nested pointer}} +// expected-error@-3 {{assigning '__local int *__local *__private' to '__generic int *__generic *__private' changes address space of nested pointer}} #endif #endif var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int; diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl --- a/clang/test/SemaOpenCL/address-spaces.cl +++ b/clang/test/SemaOpenCL/address-spaces.cl @@ -10,15 +10,15 @@ int *ip; #if ((!__OPENCL_CPP_VERSION__) && (__OPENCL_C_VERSION__ < 200)) - ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}} - ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}} - ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}} + ip = gip; // expected-error {{assigning '__global int *__private' to '__private int *__private' changes address space of pointer}} + ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}} + ip = &ci; // expected-error {{assigning '__constant int *' to '__private int *__private' changes address space of pointer}} #else ip = gip; ip = &li; ip = &ci; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}} +// expected-error@-2 {{assigning '__constant int *' to '__generic int *__private' changes address space of pointer}} #else // expected-error@-4 {{assigning to '__generic int *' from incompatible type '__constant int *'}} #endif @@ -46,9 +46,9 @@ #endif g = (__global int *)p; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting 'int *' to type '__global int *' changes address space of pointer}} +// expected-error@-2 {{casting '__private int *' to type '__global int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}} #endif l = (__local int *)g; #if !__OPENCL_CPP_VERSION__ @@ -70,9 +70,9 @@ #endif l = (__local int *)p; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting 'int *' to type '__local int *' changes address space of pointer}} +// expected-error@-2 {{casting '__private int *' to type '__local int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__private int *' to '__local int *' converts between mismatching address spaces}} #endif c = (__constant int *)g; #if !__OPENCL_CPP_VERSION__ @@ -88,33 +88,33 @@ #endif c = (__constant int *)p; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting 'int *' to type '__constant int *' changes address space of pointer}} +// expected-error@-2 {{casting '__private int *' to type '__constant int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__private int *' to '__constant int *' converts between mismatching address spaces}} #endif p = (__private int *)g; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting '__global int *' to type 'int *' changes address space of pointer}} +// expected-error@-2 {{casting '__global int *' to type '__private int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__global int *' to '__private int *' converts between mismatching address spaces}} #endif p = (__private int *)l; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting '__local int *' to type 'int *' changes address space of pointer}} +// expected-error@-2 {{casting '__local int *' to type '__private int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__local int *' to '__private int *' converts between mismatching address spaces}} #endif p = (__private int *)c; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting '__constant int *' to type 'int *' changes address space of pointer}} +// expected-error@-2 {{casting '__constant int *' to type '__private int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from '__constant int *' to '__private int *' converts between mismatching address spaces}} #endif p = (__private int *)cc; #if !__OPENCL_CPP_VERSION__ -// expected-error@-2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}} +// expected-error@-2 {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}} #else -// expected-error@-4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}} +// expected-error@-4 {{C-style cast from 'const __constant int *' to '__private int *' converts between mismatching address spaces}} #endif } @@ -126,98 +126,98 @@ #if !__OPENCL_CPP_VERSION__ void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { - g = gg; // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}} - g = l; // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}} - g = ll; // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}} - g = gg_f; // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}} - g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}} + g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *__private' changes address space of pointer}} + g = l; // expected-error {{assigning '__local int *__private' to '__global int *__private' changes address space of pointer}} + g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private' changes address space of pointer}} + g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *__private' changes address space of pointer}} + g = (__global int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__global int *' changes address space of pointer}} - gg = g; // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}} - gg = l; // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}} - gg = ll; // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}} - gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}} + gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *__private' changes address space of pointer}} + gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *__private' changes address space of pointer}} + gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *__private' changes address space of nested pointer}} + gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int *__private *__private' from '__global float *__private *__private'}} gg = (__global int * __private *)gg_f; - l = g; // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}} - l = gg; // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}} - l = ll; // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}} - l = gg_f; // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}} - l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}} + l = g; // expected-error {{assigning '__global int *__private' to '__local int *__private' changes address space of pointer}} + l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private' changes address space of pointer}} + l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *__private' changes address space of pointer}} + l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private' changes address space of pointer}} + l = (__local int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__local int *' changes address space of pointer}} - ll = g; // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}} - ll = gg; // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}} - ll = l; // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}} - ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}} - ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float **' to type '__local int **' discards qualifiers in nested pointer types}} + ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *__private' changes address space of pointer}} + ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}} + ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *__private' changes address space of pointer}} + ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}} + ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float *__private *' to type '__local int *__private *' discards qualifiers in nested pointer types}} - gg_f = g; // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}} - gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}} - gg_f = l; // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}} - gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}} + gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *__private' changes address space of pointer}} + gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float *__private *__private' from '__global int *__private *__private'}} + gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *__private' changes address space of pointer}} + gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *__private' changes address space of nested pointer}} gg_f = (__global float * __private *)gg; // FIXME: This doesn't seem right. This should be an error, not a warning. __local int * __global * __private * lll; - lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}} + lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global *__private *__private' from '__global int *__private *__private'}} typedef __local int * l_t; typedef __global int * g_t; __private l_t * pl; __private g_t * pg; - gg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to '__global int **' changes address space of nested pointer}} - pl = gg; // expected-error {{assigning '__global int **' to 'l_t *' (aka '__local int **') changes address space of nested pointer}} + gg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *__private' changes address space of nested pointer}} + pl = gg; // expected-error {{assigning '__global int *__private *__private' to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}} gg = pg; pg = gg; - pg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to 'g_t *' (aka '__global int **') changes address space of nested pointer}} - pl = pg; // expected-error {{assigning 'g_t *' (aka '__global int **') to 'l_t *' (aka '__local int **') changes address space of nested pointer}} + pg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__private g_t *__private' (aka '__global int *__private *__private') changes address space of nested pointer}} + pl = pg; // expected-error {{assigning '__private g_t *__private' (aka '__global int *__private *__private') to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}} ll = (__local int * __private *)(void *)gg; void *vp = ll; } #else void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { - g = gg; // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}} - g = l; // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}} - g = ll; // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}} - g = gg_f; // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}} - g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__global int *' converts between mismatching address spaces}} + g = gg; // expected-error {{assigning to '__global int *' from incompatible type '__global int *__private *__private'}} + g = l; // expected-error {{assigning to '__global int *' from incompatible type '__local int *__private'}} + g = ll; // expected-error {{assigning to '__global int *' from incompatible type '__local int *__private *__private'}} + g = gg_f; // expected-error {{assigning to '__global int *' from incompatible type '__global float *__private *__private'}} + g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__global int *' converts between mismatching address spaces}} - gg = g; // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}} - gg = l; // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}} - gg = ll; // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}} - gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}} + gg = g; // expected-error {{assigning to '__global int *__private *' from incompatible type '__global int *__private'; take the address with &}} + gg = l; // expected-error {{assigning to '__global int *__private *' from incompatible type '__local int *__private'}} + gg = ll; // expected-error {{assigning to '__global int *__private *' from incompatible type '__local int *__private *__private'}} + gg = gg_f; // expected-error {{assigning to '__global int *__private *' from incompatible type '__global float *__private *__private'}} gg = (__global int * __private *)gg_f; - l = g; // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}} - l = gg; // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}} - l = ll; // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}} - l = gg_f; // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}} - l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__local int *' converts between mismatching address spaces}} + l = g; // expected-error {{assigning to '__local int *' from incompatible type '__global int *__private'}} + l = gg; // expected-error {{assigning to '__local int *' from incompatible type '__global int *__private *__private'}} + l = ll; // expected-error {{assigning to '__local int *' from incompatible type '__local int *__private *__private'}} + l = gg_f; // expected-error {{assigning to '__local int *' from incompatible type '__global float *__private *__private'}} + l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}} - ll = g; // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}} - ll = gg; // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}} - ll = l; // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}} - ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}} + ll = g; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private'}} + ll = gg; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private *__private'}} + ll = l; // expected-error {{assigning to '__local int *__private *' from incompatible type '__local int *__private'; take the address with &}} + ll = gg_f; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global float *__private *__private'}} // FIXME: The below becomes a reinterpret_cast, and therefore does not emit an error // even though the address space mismatches in the nested pointers. ll = (__local int * __private *)gg; - gg_f = g; // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}} - gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}} - gg_f = l; // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}} - gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}} + gg_f = g; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private'}} + gg_f = gg; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private *__private'}} + gg_f = l; // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private'}} + gg_f = ll; // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private *__private'}} gg_f = (__global float * __private *)gg; typedef __local int * l_t; typedef __global int * g_t; __private l_t * pl; __private g_t * pg; - gg = pl; // expected-error {{assigning to '__global int **' from incompatible type 'l_t *' (aka '__local int **')}} - pl = gg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type '__global int **'}} + gg = pl; // expected-error {{assigning to '__global int *__private *' from incompatible type '__private l_t *__private' (aka '__local int *__private *__private')}} + pl = gg; // expected-error {{assigning to '__private l_t *' (aka '__local int *__private *') from incompatible type '__global int *__private *__private'}} gg = pg; pg = gg; - pg = pl; // expected-error {{assigning to 'g_t *' (aka '__global int **') from incompatible type 'l_t *' (aka '__local int **')}} - pl = pg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type 'g_t *' (aka '__global int **')}} + pg = pl; // expected-error {{assigning to '__private g_t *' (aka '__global int *__private *') from incompatible type '__private l_t *__private' (aka '__local int *__private *__private')}} + pl = pg; // expected-error {{assigning to '__private l_t *' (aka '__local int *__private *') from incompatible type '__private g_t *__private' (aka '__global int *__private *__private')}} ll = (__local int * __private *)(void *)gg; void *vp = ll; diff --git a/clang/test/SemaOpenCL/arithmetic-conversions.cl b/clang/test/SemaOpenCL/arithmetic-conversions.cl --- a/clang/test/SemaOpenCL/arithmetic-conversions.cl +++ b/clang/test/SemaOpenCL/arithmetic-conversions.cl @@ -14,10 +14,10 @@ kernel void foo5(float2 in, global float2 *out) { float* f; - *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('float *' and 'float2' (vector of 2 'float' values))}} + *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private float *' and 'float2' (vector of 2 'float' values))}} } kernel void foo6(int2 in, global int2 *out) { int* f; - *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('int *' and 'int2' (vector of 2 'int' values))}} + *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private int *' and 'int2' (vector of 2 'int' values))}} } diff --git a/clang/test/SemaOpenCL/as_type.cl b/clang/test/SemaOpenCL/as_type.cl --- a/clang/test/SemaOpenCL/as_type.cl +++ b/clang/test/SemaOpenCL/as_type.cl @@ -1,14 +1,14 @@ // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -finclude-default-header -o - -verify -fsyntax-only char3 f1(char16 x) { - return __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and 'char16' (vector of 16 'char' values) must match}} + return __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and '__private char16' (vector of 16 'char' values) must match}} } char16 f3(int x) { - return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and 'int' must match}} + return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and '__private int' must match}} } void foo() { char src = 1; - int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and 'char' must match}} + int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and '__private char' must match}} } diff --git a/clang/test/SemaOpenCL/atomic-ops.cl b/clang/test/SemaOpenCL/atomic-ops.cl --- a/clang/test/SemaOpenCL/atomic-ops.cl +++ b/clang/test/SemaOpenCL/atomic-ops.cl @@ -82,13 +82,13 @@ bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); - bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}} - (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}} + bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}} + (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}} bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); - bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}} - (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}} + bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}} + (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}} // Pointers to different address spaces are allowed. bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); @@ -107,7 +107,7 @@ float forder; (void)__opencl_atomic_load(Ap, forder, memory_scope_work_group); struct S s; - (void)__opencl_atomic_load(Ap, s, memory_scope_work_group); // expected-error {{passing 'struct S' to parameter of incompatible type 'int'}} + (void)__opencl_atomic_load(Ap, s, memory_scope_work_group); // expected-error {{passing '__private struct S' to parameter of incompatible type 'int'}} (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_work_group); (void)__opencl_atomic_load(Ap, memory_order_acquire, memory_scope_work_group); @@ -187,7 +187,7 @@ (void)__opencl_atomic_load(Ap, memory_order_relaxed, 1.0f); (void)__opencl_atomic_load(Ap, memory_order_relaxed, fscope); struct S s; - (void)__opencl_atomic_load(Ap, memory_order_relaxed, s); //expected-error{{passing 'struct S' to parameter of incompatible type 'int'}} + (void)__opencl_atomic_load(Ap, memory_order_relaxed, s); //expected-error{{passing '__private struct S' to parameter of incompatible type 'int'}} } void nullPointerWarning(atomic_int *Ap, int *p, int val) { diff --git a/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl b/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl --- a/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl +++ b/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl @@ -91,7 +91,7 @@ }, c, 1024L); #ifdef WCONV -// expected-warning-re@-2{{implicit conversion changes signedness: 'char' to 'unsigned {{int|long}}'}} +// expected-warning-re@-2{{implicit conversion changes signedness: '__private char' to 'unsigned {{int|long}}'}} #endif #define UINT_MAX 4294967295 diff --git a/clang/test/SemaOpenCL/clk_event_t.cl b/clang/test/SemaOpenCL/clk_event_t.cl --- a/clang/test/SemaOpenCL/clk_event_t.cl +++ b/clang/test/SemaOpenCL/clk_event_t.cl @@ -12,7 +12,7 @@ clk_event_t ce2; clk_event_t ce3 = CLK_NULL_EVENT; - if (e == ce1) { // expected-error {{invalid operands to binary expression ('event_t' and 'clk_event_t')}} + if (e == ce1) { // expected-error {{invalid operands to binary expression ('__private event_t' and '__private clk_event_t')}} return 9; } diff --git a/clang/test/SemaOpenCL/event_t.cl b/clang/test/SemaOpenCL/event_t.cl --- a/clang/test/SemaOpenCL/event_t.cl +++ b/clang/test/SemaOpenCL/event_t.cl @@ -1,6 +1,6 @@ // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}} +event_t glb_evt; // expected-error {{the '__private event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}} constant struct evt_s { event_t evt; // expected-error {{the 'event_t' type cannot be used to declare a structure or union field}} @@ -8,7 +8,7 @@ void foo(event_t evt); // expected-note {{passing argument to parameter 'evt' here}} -void kernel ker(event_t argevt) { // expected-error {{'event_t' cannot be used as the type of a kernel parameter}} +void kernel ker(event_t argevt) { // expected-error {{'__private event_t' cannot be used as the type of a kernel parameter}} event_t e; constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}} expected-error{{variable in constant address space must be initialized}} foo(e); diff --git a/clang/test/SemaOpenCL/extension-begin.cl b/clang/test/SemaOpenCL/extension-begin.cl --- a/clang/test/SemaOpenCL/extension-begin.cl +++ b/clang/test/SemaOpenCL/extension-begin.cl @@ -44,7 +44,7 @@ struct A test_A2; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}} const struct A test_A_local; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}} TypedefOfA test_typedef_A; // expected-error {{use of type 'TypedefOfA' (aka 'struct A') requires my_ext extension to be enabled}} - PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const struct A *') requires my_ext extension to be enabled}} + PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const __private struct A *') requires my_ext extension to be enabled}} f(); // expected-error {{use of declaration 'f' requires my_ext extension to be enabled}} g(0); // expected-error {{no matching function for call to 'g'}} // expected-note@extension-begin.h:18 {{candidate unavailable as it requires OpenCL extension 'my_ext' to be enabled}} diff --git a/clang/test/SemaOpenCL/half.cl b/clang/test/SemaOpenCL/half.cl --- a/clang/test/SemaOpenCL/half.cl +++ b/clang/test/SemaOpenCL/half.cl @@ -4,13 +4,13 @@ constant float f = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}} half half_disabled(half *p, // expected-error{{declaring function return value of type 'half' is not allowed}} - half h) // expected-error{{declaring function parameter of type 'half' is not allowed}} + half h) // expected-error{{declaring function parameter of type '__private half' is not allowed}} { - half a[2]; // expected-error{{declaring variable of type 'half [2]' is not allowed}} - half b; // expected-error{{declaring variable of type 'half' is not allowed}} - *p; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}} + half a[2]; // expected-error{{declaring variable of type '__private half [2]' is not allowed}} + half b; // expected-error{{declaring variable of type '__private half' is not allowed}} + *p; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}} *p = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}} - p[1]; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}} + p[1]; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}} p[1] = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}} float c = 1.0f; @@ -26,7 +26,7 @@ } kernel void half_disabled_kernel(global half *p, - half h); // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}} + half h); // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}} // Exactly the same as above but with the cl_khr_fp16 extension enabled. #pragma OPENCL EXTENSION cl_khr_fp16 : enable diff --git a/clang/test/SemaOpenCL/images.cl b/clang/test/SemaOpenCL/images.cl --- a/clang/test/SemaOpenCL/images.cl +++ b/clang/test/SemaOpenCL/images.cl @@ -7,26 +7,26 @@ void imgage_access_test(image2d_t img2dro, image3d_t img3dro) { img2d_ro(img2dro); // read_only = read_only - img2d_ro(img3dro); // expected-error{{passing '__read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}} + img2d_ro(img3dro); // expected-error{{passing '__private __read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}} } kernel void read_only_access_test(read_only image2d_t img) { img2d_ro(img); // read_only = read_only - img2d_wo(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}} - img2d_rw(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}} + img2d_wo(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}} + img2d_rw(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}} img2d_default(img); // read_only = read_only } kernel void write_only_access_test(write_only image2d_t img) { - img2d_ro(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}} + img2d_ro(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}} img2d_wo(img); // write_only = write_only - img2d_rw(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}} - img2d_default(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}} + img2d_rw(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}} + img2d_default(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}} } kernel void read_write_access_test(read_write image2d_t img) { - img2d_ro(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}} - img2d_wo(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}} + img2d_ro(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}} + img2d_wo(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}} img2d_rw(img); //read_write = read_write - img2d_default(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}} + img2d_default(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}} } diff --git a/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl b/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl --- a/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl +++ b/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl @@ -16,34 +16,34 @@ // negative test cases for initializers void foo(char c, float f, void* v, struct st ss) { intel_sub_group_avc_mce_payload_t payload_mce = 0; // No zero initializer for mce types - // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}} intel_sub_group_avc_ime_payload_t payload_ime = 1; // No literal initializer for *payload_t types - // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}} intel_sub_group_avc_ref_payload_t payload_ref = f; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'float'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_payload_t' with an expression of incompatible type '__private float'}} intel_sub_group_avc_sic_payload_t payload_sic = ss; - // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_payload_t' with an expression of incompatible type 'struct st'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_sic_payload_t' with an expression of incompatible type '__private struct st'}} intel_sub_group_avc_mce_result_t result_mce = 0; // No zero initializer for mce types - // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}} intel_sub_group_avc_ime_result_t result_ime = 1; // No literal initializer for *result_t types - // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}} intel_sub_group_avc_ref_result_t result_ref = f; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_result_t' with an expression of incompatible type 'float'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_result_t' with an expression of incompatible type '__private float'}} intel_sub_group_avc_sic_result_t result_sic = ss; - // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_result_t' with an expression of incompatible type 'struct st'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_sic_result_t' with an expression of incompatible type '__private struct st'}} intel_sub_group_avc_ime_result_single_reference_streamout_t sstreamout = v; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type 'void *'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type '__private void *__private'}} intel_sub_group_avc_ime_result_dual_reference_streamout_t dstreamin_list = {0x0, 0x1}; // expected-warning@-1 {{excess elements in struct initializer}} intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list2 = {}; // expected-error@-1 {{scalar initializer cannot be empty}} intel_sub_group_avc_ime_single_reference_streamin_t dstreamin_list3 = {c}; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type 'char'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type '__private char'}} intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list4 = {1}; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}} } // negative tests for initializers and assignment @@ -53,12 +53,12 @@ intel_sub_group_avc_ime_payload_t payload_ime; intel_sub_group_avc_ref_payload_t payload_ref = payload_ime; - // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'intel_sub_group_avc_ime_payload_t'}} + // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_payload_t' with an expression of incompatible type '__private intel_sub_group_avc_ime_payload_t'}} intel_sub_group_avc_sic_result_t result_sic; intel_sub_group_avc_ime_result_t result_ime; result_sic = result_ime; - // expected-error@-1 {{assigning to 'intel_sub_group_avc_sic_result_t' from incompatible type 'intel_sub_group_avc_ime_result_t'}} + // expected-error@-1 {{assigning to '__private intel_sub_group_avc_sic_result_t' from incompatible type '__private intel_sub_group_avc_ime_result_t'}} } // Using 0x0 directly allows us not to include opencl-c.h header and not to diff --git a/clang/test/SemaOpenCL/invalid-block.cl b/clang/test/SemaOpenCL/invalid-block.cl --- a/clang/test/SemaOpenCL/invalid-block.cl +++ b/clang/test/SemaOpenCL/invalid-block.cl @@ -12,7 +12,7 @@ }; f0(bl1); f0(bl2); - bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (__generic ^const)(void)' and 'int (__generic ^const)(void)')}} + bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (__generic ^const __private)(void)' and 'int (__generic ^const __private)(void)')}} int (^const bl3)(); // expected-error{{invalid block variable declaration - must be initialized}} } @@ -53,18 +53,18 @@ bl2_t bl2 = ^(int i) { return 2; }; - bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(int)') type is invalid in OpenCL}} + bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(__private int)') type is invalid in OpenCL}} int tmp = i ? bl1(i) // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}} : bl2(i); // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}} } // A block pointer type and all pointer operations are disallowed -void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}} +void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (__generic ^const)(__private int)') is invalid in OpenCL}} bl2_t bl = ^(int i) { return 1; }; - bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}} - *bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}} - &bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}} + bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (__generic ^const)(__private int)') is invalid in OpenCL}} + *bl; // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private int)') to unary expression}} + &bl; // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private int)') to unary expression}} } // A block can't reference another block kernel void f7() { diff --git a/clang/test/SemaOpenCL/invalid-image.cl b/clang/test/SemaOpenCL/invalid-image.cl --- a/clang/test/SemaOpenCL/invalid-image.cl +++ b/clang/test/SemaOpenCL/invalid-image.cl @@ -5,12 +5,12 @@ void test1(image1d_t *i) {} // expected-error-re{{pointer to type '{{__generic __read_only|__read_only}} image1d_t' is invalid in OpenCL}} void test2(image1d_t i) { - image1d_t ti; // expected-error{{type '__read_only image1d_t' can only be used as a function parameter}} + image1d_t ti; // expected-error{{type '__private __read_only image1d_t' can only be used as a function parameter}} image1d_t ai[] = {i, i}; // expected-error{{array of '__read_only image1d_t' type is invalid in OpenCL}} - i=i; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and '__read_only image1d_t')}} - i+1; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and 'int')}} - &i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}} - *i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}} + i=i; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and '__private __read_only image1d_t')}} + i+1; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and 'int')}} + &i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}} + *i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}} } image1d_t test3() {} // expected-error{{declaring function return value of type '__read_only image1d_t' is not allowed}} diff --git a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl --- a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl +++ b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown -kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}} +kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}} #pragma OPENCL EXTENSION cl_khr_fp16 : enable @@ -9,37 +9,37 @@ // bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t // or a struct / union with any of these types in them -typedef __SIZE_TYPE__ size_t; // expected-note{{'size_t' (aka 'unsigned int') declared here}} +typedef __SIZE_TYPE__ size_t; // expected-note{{'__private size_t' (aka '__private unsigned int') declared here}} // expected-note@-1{{'size_t' (aka 'unsigned int') declared here}} -typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'ptrdiff_t' (aka 'int') declared here}} -typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'intptr_t' (aka 'int') declared here}} -typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'uintptr_t' (aka 'unsigned int') declared here}} +typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'__private ptrdiff_t' (aka '__private int') declared here}} +typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'__private intptr_t' (aka '__private int') declared here}} +typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'__private uintptr_t' (aka '__private unsigned int') declared here}} -kernel void size_t_arg(size_t x) {} // expected-error{{'size_t' (aka 'unsigned int') cannot be used as the type of a kernel parameter}} +kernel void size_t_arg(size_t x) {} // expected-error{{'__private size_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}} -kernel void ptrdiff_t_arg(ptrdiff_t x) {} // expected-error{{'ptrdiff_t' (aka 'int') cannot be used as the type of a kernel parameter}} +kernel void ptrdiff_t_arg(ptrdiff_t x) {} // expected-error{{'__private ptrdiff_t' (aka '__private int') cannot be used as the type of a kernel parameter}} -kernel void intptr_t_arg(intptr_t x) {} // expected-error{{'intptr_t' (aka 'int') cannot be used as the type of a kernel parameter}} +kernel void intptr_t_arg(intptr_t x) {} // expected-error{{'__private intptr_t' (aka '__private int') cannot be used as the type of a kernel parameter}} -kernel void uintptr_t_arg(uintptr_t x) {} // expected-error{{'uintptr_t' (aka 'unsigned int') cannot be used as the type of a kernel parameter}} +kernel void uintptr_t_arg(uintptr_t x) {} // expected-error{{'__private uintptr_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}} typedef size_t size_ty; struct SizeTStruct { // expected-note{{within field of type 'SizeTStruct' declared here}} size_ty s; // expected-note{{field of illegal type 'size_ty' (aka 'unsigned int') declared here}} }; -kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'struct SizeTStruct' cannot be used as the type of a kernel parameter}} +kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'__private struct SizeTStruct' cannot be used as the type of a kernel parameter}} union SizeTUnion { // expected-note{{within field of type 'SizeTUnion' declared here}} size_t s; // expected-note{{field of illegal type 'size_t' (aka 'unsigned int') declared here}} float f; }; -kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'union SizeTUnion' cannot be used as the type of a kernel parameter}} +kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'__private union SizeTUnion' cannot be used as the type of a kernel parameter}} typedef size_t s_ty; // expected-note{{'s_ty' (aka 'unsigned int') declared here}} -typedef s_ty ss_ty; // expected-note{{'ss_ty' (aka 'unsigned int') declared here}} -kernel void typedef_to_size_t(ss_ty s) {} // expected-error{{'ss_ty' (aka 'unsigned int') cannot be used as the type of a kernel parameter}} +typedef s_ty ss_ty; // expected-note{{'__private ss_ty' (aka '__private unsigned int') declared here}} +kernel void typedef_to_size_t(ss_ty s) {} // expected-error{{'__private ss_ty' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}} -kernel void bool_arg(bool x) { } // expected-error{{'bool' cannot be used as the type of a kernel parameter}} +kernel void bool_arg(bool x) { } // expected-error{{'__private bool' cannot be used as the type of a kernel parameter}} // half kernel argument is allowed when cl_khr_fp16 is enabled. kernel void half_arg(half x) { } @@ -49,7 +49,7 @@ bool x; // expected-note{{field of illegal type 'bool' declared here}} } ContainsBool; -kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'ContainsBool' (aka 'struct ContainsBool') cannot be used as the type of a kernel parameter}} +kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'__private ContainsBool' (aka '__private struct ContainsBool') cannot be used as the type of a kernel parameter}} @@ -65,14 +65,14 @@ typedef struct Foo // expected-note{{within field of type 'Foo' declared here}} { - int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}} + int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}} } Foo; kernel void pointer_in_struct_arg(Foo arg) { } // expected-error{{struct kernel parameters may not contain pointers}} typedef union FooUnion // expected-note{{within field of type 'FooUnion' declared here}} { - int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}} + int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}} } FooUnion; kernel void pointer_in_union_arg(FooUnion arg) { }// expected-error{{union kernel parameters may not contain pointers}} @@ -82,7 +82,7 @@ int x; struct InnerNestedPointer { - int* ptrField; // expected-note 3 {{field of illegal pointer type 'int *' declared here}} + int* ptrField; // expected-note 3 {{field of illegal pointer type '__private int *' declared here}} } inner; // expected-note 3 {{within field of type 'struct InnerNestedPointer' declared here}} } NestedPointer; @@ -96,7 +96,7 @@ struct InnerNestedPointerComplex { int innerFoo; - int* innerPtrField; // expected-note{{field of illegal pointer type 'int *' declared here}} + int* innerPtrField; // expected-note{{field of illegal pointer type '__private int *' declared here}} } inner; // expected-note{{within field of type 'struct InnerNestedPointerComplex' declared here}} float y; @@ -114,10 +114,10 @@ } inner; // expected-note 2 {{within field of type 'struct InnerNestedBool' declared here}} } NestedBool; -kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}} +kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}} // Warning emitted again for argument used in other kernel -kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}} +kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}} // Check for note with a struct not defined inside the struct @@ -132,7 +132,7 @@ NestedBool2Inner inner; // expected-note{{within field of type 'NestedBool2Inner' (aka 'struct NestedBool2Inner') declared here}} } NestedBool2; -kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'NestedBool2' (aka 'struct NestedBool2') cannot be used as the type of a kernel parameter}} +kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'__private NestedBool2' (aka '__private struct NestedBool2') cannot be used as the type of a kernel parameter}} struct InnerInner @@ -167,8 +167,8 @@ struct ArrayOfPtr // expected-note{{within field of type 'ArrayOfPtr' declared here}} { - float *arr[3]; // expected-note{{field of illegal type 'float *[3]' declared here}} - // expected-note@-1{{field of illegal type 'float *[3]' declared here}} + float *arr[3]; // expected-note{{field of illegal type '__private float *[3]' declared here}} + // expected-note@-1{{field of illegal type '__private float *[3]' declared here}} }; kernel void array_of_ptr(struct ArrayOfPtr arr) {} // expected-error{{struct kernel parameters may not contain pointers}} diff --git a/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl b/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl --- a/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl +++ b/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl @@ -12,22 +12,22 @@ read_pipe(tmp, p); // expected-error {{first argument to 'read_pipe' must be a pipe type}} read_pipe(p); // expected-error {{invalid number of arguments to function: 'read_pipe'}} read_pipe(p, rid, tmp, ptr); - read_pipe(p, tmp, tmp, ptr); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having 'int')}} - read_pipe(p, rid, rid, ptr); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having 'reserve_id_t')}} - read_pipe(p, tmp); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having 'int')}} + read_pipe(p, tmp, tmp, ptr); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having '__private int')}} + read_pipe(p, rid, rid, ptr); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}} + read_pipe(p, tmp); // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having '__private int')}} write_pipe(p, ptr); // expected-error {{invalid pipe access modifier (expecting write_only)}} write_pipe(p, rid, tmp, ptr); // expected-error {{invalid pipe access modifier (expecting write_only)}} // reserve_read/write_pipe reserve_read_pipe(p, tmp); - reserve_read_pipe(p, ptr); // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *')}} + reserve_read_pipe(p, ptr); // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *__private')}} work_group_reserve_read_pipe(tmp, tmp); // expected-error{{first argument to 'work_group_reserve_read_pipe' must be a pipe type}} sub_group_reserve_write_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting write_only)}} // commit_read/write_pipe commit_read_pipe(p, rid); commit_read_pipe(tmp, rid); // expected-error{{first argument to 'commit_read_pipe' must be a pipe type}} - work_group_commit_read_pipe(p, tmp); // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having 'int')}} + work_group_commit_read_pipe(p, tmp); // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having '__private int')}} sub_group_commit_write_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting write_only)}} } @@ -41,22 +41,22 @@ write_pipe(tmp, p); // expected-error {{first argument to 'write_pipe' must be a pipe type}} write_pipe(p); // expected-error {{invalid number of arguments to function: 'write_pipe'}} write_pipe(p, rid, tmp, ptr); - write_pipe(p, tmp, tmp, ptr); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having 'int')}} - write_pipe(p, rid, rid, ptr); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having 'reserve_id_t')}} - write_pipe(p, tmp); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having 'int')}} + write_pipe(p, tmp, tmp, ptr); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having '__private int')}} + write_pipe(p, rid, rid, ptr); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}} + write_pipe(p, tmp); // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having '__private int')}} read_pipe(p, ptr); // expected-error {{invalid pipe access modifier (expecting read_only)}} read_pipe(p, rid, tmp, ptr); // expected-error {{invalid pipe access modifier (expecting read_only)}} // reserve_read/write_pipe reserve_write_pipe(p, tmp); - reserve_write_pipe(p, ptr); // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *')}} + reserve_write_pipe(p, ptr); // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *__private')}} work_group_reserve_write_pipe(tmp, tmp); // expected-error{{first argument to 'work_group_reserve_write_pipe' must be a pipe type}} sub_group_reserve_read_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting read_only)}} // commit_read/write_pipe commit_write_pipe(p, rid); commit_write_pipe(tmp, rid); // expected-error{{first argument to 'commit_write_pipe' must be a pipe type}} - work_group_commit_write_pipe(p, tmp); // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having 'int')}} + work_group_commit_write_pipe(p, tmp); // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having '__private int')}} sub_group_commit_read_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting read_only)}} } diff --git a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl --- a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl +++ b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl @@ -6,7 +6,7 @@ extern pipe write_only int get_pipe(); // expected-error-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}} expected-error{{'write_only' attribute only applies to parameters and typedefs}} -kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}} +kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'__private reserve_id_t' cannot be used as the type of a kernel parameter}} } void test1(pipe int *p) {// expected-error {{pipes packet types cannot be of reference type}} @@ -16,15 +16,15 @@ void test3(int pipe p) {// expected-error {{cannot combine with previous 'int' declaration specifier}} } void test4() { - pipe int p; // expected-error {{type 'read_only pipe int' can only be used as a function parameter}} + pipe int p; // expected-error {{type '__private read_only pipe int' can only be used as a function parameter}} //TODO: fix parsing of this pipe int (*p); } void test5(pipe int p) { - p+p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}} - p=p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}} - &p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}} - *p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}} + p+p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}} + p=p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}} + &p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}} + *p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}} } typedef pipe int pipe_int_t; @@ -32,7 +32,7 @@ bool test_id_comprision(void) { reserve_id_t id1, id2; - return (id1 == id2); // expected-error {{invalid operands to binary expression ('reserve_id_t' and 'reserve_id_t')}} + return (id1 == id2); // expected-error {{invalid operands to binary expression ('__private reserve_id_t' and '__private reserve_id_t')}} } // Tests ASTContext::mergeTypes rejects this. diff --git a/clang/test/SemaOpenCL/null_literal.cl b/clang/test/SemaOpenCL/null_literal.cl --- a/clang/test/SemaOpenCL/null_literal.cl +++ b/clang/test/SemaOpenCL/null_literal.cl @@ -11,14 +11,14 @@ constant int* ptr3 = NULL; -constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}} +constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *__private' with an expression of type '__global void *' changes address space of pointer}} #ifdef CL20 // Accept explicitly pointer to generic address space in OpenCL v2.0. global int* ptr5 = (generic void*)0; #endif -global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}} +global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *__private' with an expression of type '__local void *' changes address space of pointer}} bool cmp = ptr1 == NULL; diff --git a/clang/test/SemaOpenCL/null_queue.cl b/clang/test/SemaOpenCL/null_queue.cl --- a/clang/test/SemaOpenCL/null_queue.cl +++ b/clang/test/SemaOpenCL/null_queue.cl @@ -4,13 +4,13 @@ void queue_arg(queue_t); // expected-note {{passing argument to parameter here}} void init() { - queue_t q1 = 1; // expected-error{{initializing 'queue_t' with an expression of incompatible type 'int'}} + queue_t q1 = 1; // expected-error{{initializing '__private queue_t' with an expression of incompatible type 'int'}} queue_t q = 0; } void assign() { queue_t q2, q3; - q2 = 5; // expected-error{{assigning to 'queue_t' from incompatible type 'int'}} + q2 = 5; // expected-error{{assigning to '__private queue_t' from incompatible type 'int'}} q3 = 0; q2 = q3 = 0; } @@ -21,7 +21,7 @@ get_default_queue() == 1 && // expected-error{{invalid operands to binary expression ('queue_t' and 'int')}} q4 == q5 && q4 != 0 && - q4 != 0.0f; // expected-error{{invalid operands to binary expression ('queue_t' and 'float')}} + q4 != 0.0f; // expected-error{{invalid operands to binary expression ('__private queue_t' and 'float')}} } void call() { diff --git a/clang/test/SemaOpenCL/predefined-expr.cl b/clang/test/SemaOpenCL/predefined-expr.cl --- a/clang/test/SemaOpenCL/predefined-expr.cl +++ b/clang/test/SemaOpenCL/predefined-expr.cl @@ -2,7 +2,7 @@ // RUN: %clang_cc1 %s -verify -cl-std=CL2.0 void f() { - char *f1 = __func__; //expected-error-re{{initializing '{{__generic char|char}} *' with an expression of type 'const __constant char *' changes address space of pointer}} - constant char *f2 = __func__; //expected-warning{{initializing '__constant char *' with an expression of type 'const __constant char [2]' discards qualifiers}} + char *f1 = __func__; //expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}} + constant char *f2 = __func__; //expected-warning{{initializing '__constant char *__private' with an expression of type 'const __constant char [2]' discards qualifiers}} constant const char *f3 = __func__; } diff --git a/clang/test/SemaOpenCL/queue_t_overload.cl b/clang/test/SemaOpenCL/queue_t_overload.cl --- a/clang/test/SemaOpenCL/queue_t_overload.cl +++ b/clang/test/SemaOpenCL/queue_t_overload.cl @@ -1,7 +1,7 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only -void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}} -void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}} +void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}} +void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}} void kernel ker(__local char *src1, __local float *src2, __global int *src3) { queue_t q; diff --git a/clang/test/SemaOpenCL/shifts.cl b/clang/test/SemaOpenCL/shifts.cl --- a/clang/test/SemaOpenCL/shifts.cl +++ b/clang/test/SemaOpenCL/shifts.cl @@ -45,7 +45,7 @@ // ** Negative tests ** char2 ntest01(char c, char2 s) { - return c << s; // expected-error {{requested shift is a vector of type 'char2' (vector of 2 'char' values) but the first operand is not a vector ('char')}} + return c << s; // expected-error {{requested shift is a vector of type '__private char2' (vector of 2 'char' values) but the first operand is not a vector ('__private char')}} } char3 ntest02(char3 c, char2 s) { diff --git a/clang/test/SemaOpenCL/to_addr_builtin.cl b/clang/test/SemaOpenCL/to_addr_builtin.cl --- a/clang/test/SemaOpenCL/to_addr_builtin.cl +++ b/clang/test/SemaOpenCL/to_addr_builtin.cl @@ -13,7 +13,7 @@ glob = to_global(glob, loc); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 // expected-error@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}} - // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else // expected-error@-5{{invalid number of arguments to function: 'to_global'}} #endif @@ -21,46 +21,46 @@ int x; glob = to_global(x); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else // expected-error@-4{{invalid argument x to function: 'to_global', expecting a generic pointer argument}} #endif glob = to_global(con); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else // expected-error@-4{{invalid argument con to function: 'to_global', expecting a generic pointer argument}} #endif glob = to_global(con_typedef); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else // expected-error@-4{{invalid argument con_typedef to function: 'to_global', expecting a generic pointer argument}} #endif loc = to_global(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}} #else - // expected-error@-4{{assigning '__global int *' to '__local int *' changes address space of pointer}} + // expected-error@-4{{assigning '__global int *' to '__local int *__private' changes address space of pointer}} // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}} #endif loc = to_private(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 // expected-error@-2{{implicit declaration of function 'to_private' is invalid in OpenCL}} - // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}} + // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}} #else - // expected-error@-5{{assigning 'int *' to '__local int *' changes address space of pointer}} + // expected-error@-5{{assigning '__private int *' to '__local int *__private' changes address space of pointer}} // expected-warning@-6{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}} #endif loc = to_local(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 // expected-error@-2{{implicit declaration of function 'to_local' is invalid in OpenCL}} - // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}} + // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}} // expected-note@-4{{did you mean 'to_global'}} // expected-note@13{{'to_global' declared here}} #else @@ -69,15 +69,15 @@ priv = to_global(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}} #else - // expected-error@-4{{assigning '__global int *' to 'int *' changes address space of pointer}} + // expected-error@-4{{assigning '__global int *' to '__private int *__private' changes address space of pointer}} // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}} #endif priv = to_private(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}} #else // expected-warning@-4{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}} #endif @@ -85,48 +85,48 @@ priv = to_local(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}} #else - // expected-error@-4{{assigning '__local int *' to 'int *' changes address space of pointer}} + // expected-error@-4{{assigning '__local int *' to '__private int *__private' changes address space of pointer}} // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}} #endif glob = to_global(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else // expected-warning@-4{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}} #endif glob = to_private(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else - // expected-error@-4{{assigning 'int *' to '__global int *' changes address space of pointer}} + // expected-error@-4{{assigning '__private int *' to '__global int *__private' changes address space of pointer}} // expected-warning@-5{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}} #endif glob = to_local(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}} #else - // expected-error@-4{{assigning '__local int *' to '__global int *' changes address space of pointer}} + // expected-error@-4{{assigning '__local int *' to '__global int *__private' changes address space of pointer}} // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}} #endif global char *glob_c = to_global(loc); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *' with an expression of type 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *__private' with an expression of type 'int'}} #else - // expected-warning@-4{{incompatible pointer types initializing '__global char *' with an expression of type '__global int *'}} + // expected-warning@-4{{incompatible pointer types initializing '__global char *__private' with an expression of type '__global int *'}} // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}} #endif glob_wrong_ty = to_global(glob); #if __OPENCL_C_VERSION__ < CL_VERSION_2_0 - // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global float *' from 'int'}} + // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global float *__private' from 'int'}} #else - // expected-warning@-4{{incompatible pointer types assigning to '__global float *' from '__global int *'}} + // expected-warning@-4{{incompatible pointer types assigning to '__global float *__private' from '__global int *'}} // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}} #endif diff --git a/clang/test/SemaOpenCL/vec_step.cl b/clang/test/SemaOpenCL/vec_step.cl --- a/clang/test/SemaOpenCL/vec_step.cl +++ b/clang/test/SemaOpenCL/vec_step.cl @@ -26,7 +26,7 @@ int res11[vec_step(int16) == 16 ? 1 : -1]; int res12[vec_step(void) == 1 ? 1 : -1]; - int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, 'struct S' invalid}} - int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, 'int16 *' invalid}} + int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private struct S' invalid}} + int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private int16 *' invalid}} int res15 = vec_step(void(void)); // expected-error {{'vec_step' requires built-in scalar or vector type, 'void (void)' invalid}} } diff --git a/clang/test/SemaOpenCL/vector_conv_invalid.cl b/clang/test/SemaOpenCL/vector_conv_invalid.cl --- a/clang/test/SemaOpenCL/vector_conv_invalid.cl +++ b/clang/test/SemaOpenCL/vector_conv_invalid.cl @@ -7,7 +7,7 @@ void vector_conv_invalid(const global int4 *const_global_ptr) { uint4 u = (uint4)(1); - int4 i = u; // expected-error{{initializing 'int4' (vector of 4 'int' values) with an expression of incompatible type 'uint4' (vector of 4 'unsigned int' values)}} + int4 i = u; // expected-error{{initializing '__private int4' (vector of 4 'int' values) with an expression of incompatible type '__private uint4' (vector of 4 'unsigned int' values)}} int4 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' (vector of 4 'int' values) and 'uint4' (vector of 4 'unsigned int' values)}} uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' (vector of 3 'unsigned int' values) and 'uint4' (vector of 4 'unsigned int' values)}} @@ -16,7 +16,7 @@ e = (constant int4)i; e = (private int4)i; - private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const int4 *' changes address space of pointer}} - global int4 *global_ptr = const_global_ptr; // expected-warning {{initializing '__global int4 *' with an expression of type 'const __global int4 *' discards qualifiers}} + private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}} + global int4 *global_ptr = const_global_ptr; // expected-warning {{initializing '__global int4 *__private' with an expression of type 'const __global int4 *__private' discards qualifiers}} global_ptr = (global int4 *)const_global_ptr; } diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl --- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -30,8 +30,8 @@ template struct x1 { -//CHECK: -CXXMethodDecl {{.*}} operator= 'x1 &(const x1 &){{( __attribute__.*)?}} __generic' -//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1 &(const __generic x1 &){{( __attribute__.*)?}} __generic' +//CHECK: -CXXMethodDecl {{.*}} operator= 'x1 &(const x1 &__private){{( __attribute__.*)?}} __generic' +//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1 &(const __generic x1 &__private){{( __attribute__.*)?}} __generic' x1& operator=(const x1& xx) { y = xx.y; return *this; @@ -41,8 +41,8 @@ template struct x2 { -//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1 *){{( __attribute__.*)?}} __generic' -//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1 *){{( __attribute__.*)?}} __generic' +//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1 *__private){{( __attribute__.*)?}} __generic' +//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1 *__private){{( __attribute__.*)?}} __generic' void foo(x1* xx) { m[0] = *xx; } @@ -57,10 +57,10 @@ template class x3 : public T { public: - //CHECK: -CXXConstructorDecl {{.*}} x3 'void (const x3 &){{( __attribute__.*)?}} __generic' + //CHECK: -CXXConstructorDecl {{.*}} x3 'void (const x3 &__private){{( __attribute__.*)?}} __generic' x3(const x3 &t); }; -//CHECK: -CXXConstructorDecl {{.*}} x3 'void (const x3 &){{( __attribute__.*)?}} __generic' +//CHECK: -CXXConstructorDecl {{.*}} x3 'void (const x3 &__private){{( __attribute__.*)?}} __generic' template x3::x3(const x3 &t) {} @@ -68,7 +68,7 @@ T xxx(T *in1, T in2) { // This pointer can't be deduced to generic because addr space // will be taken from the template argument. - //CHECK: `-VarDecl {{.*}} i 'T *' cinit + //CHECK: `-VarDecl {{.*}} '__private T *__private' cinit T *i = in1; T ii; __private T *ptr = ⅈ @@ -92,15 +92,15 @@ } // Addr space for pointer/reference to an array -//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &)[2])' +//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &__private)[2])' void t1(const float (&fYZ)[2]); -//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *)[2])' +//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *__private)[2])' void t2(const float (*fYZ)[2]); -//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *)))[2])' +//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *__private)))[2])' void t3(float(((*fYZ)))[2]); -//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *)))[2])' +//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *__private)))[2])' void t4(float(((**fYZ)))[2]); -//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *))[2])' +//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *__private))[2])' void t5(float (*(*fYZ))[2]); __kernel void k() { diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl --- a/clang/test/SemaOpenCLCXX/address-space-lambda.cl +++ b/clang/test/SemaOpenCLCXX/address-space-lambda.cl @@ -1,6 +1,6 @@ //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s -//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic' +//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (__private int) const __generic' auto glambda = [](auto a) { return a; }; __kernel void test() { @@ -25,16 +25,16 @@ } __kernel void test_qual() { -//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const' +//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __private' auto priv1 = []() __private {}; priv1(); //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic' auto priv2 = []() __generic {}; priv2(); - auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}} + auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}} priv3(); //expected-error{{no matching function for call to object of type}} - __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in default address space}} //expected-note{{conversion candidate of type 'void (*)()'}} + __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}} const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}} __constant auto const2 = []() __generic{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}} //expected-note{{conversion candidate of type 'void (*)()'}} const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}} @@ -43,7 +43,7 @@ const3(); [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}} - [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in default address space}} + [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__private'}} [&] __private {} (); //expected-error{{lambda requires '()' before attribute specifier}} expected-error{{expected body of lambda expression}} diff --git a/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl --- a/clang/test/SemaOpenCLCXX/address-space-templates.cl +++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl @@ -19,7 +19,7 @@ template void foo3() { - __private T ii; // expected-error{{conflicting address space qualifiers are provided between types 'T' and '__global int'}} + __private T ii; // expected-error{{conflicting address space qualifiers are provided between types '__private T' and '__global int'}} } void bar() { diff --git a/clang/test/SemaOpenCLCXX/addrspace-auto.cl b/clang/test/SemaOpenCLCXX/addrspace-auto.cl --- a/clang/test/SemaOpenCLCXX/addrspace-auto.cl +++ b/clang/test/SemaOpenCLCXX/addrspace-auto.cl @@ -6,30 +6,30 @@ kernel void test() { int i; - //CHECK: VarDecl {{.*}} ai 'int':'int' + //CHECK: VarDecl {{.*}} ai '__private int':'__private int' auto ai = i; constexpr int c = 1; //CHECK: VarDecl {{.*}} used cai '__constant int':'__constant int' __constant auto cai = c; - //CHECK: VarDecl {{.*}} aii 'int':'int' + //CHECK: VarDecl {{.*}} aii '__private int':'__private int' auto aii = cai; - //CHECK: VarDecl {{.*}} ref 'int &' + //CHECK: VarDecl {{.*}} ref '__private int &__private' auto &ref = i; - //CHECK: VarDecl {{.*}} ptr 'int *' + //CHECK: VarDecl {{.*}} ptr '__private int *__private' auto *ptr = &i; - //CHECK: VarDecl {{.*}} ref_c '__constant int &' + //CHECK: VarDecl {{.*}} ref_c '__constant int &__private' auto &ref_c = cai; - //CHECK: VarDecl {{.*}} ptrptr 'int *__generic *' + //CHECK: VarDecl {{.*}} ptrptr '__private int *__generic *__private' auto **ptrptr = &ptr; - //CHECK: VarDecl {{.*}} refptr 'int *__generic &' + //CHECK: VarDecl {{.*}} refptr '__private int *__generic &__private' auto *&refptr = ptr; - //CHECK: VarDecl {{.*}} invalid gref '__global auto &' - __global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &' has incompatible initializer of type 'int'}} + //CHECK: VarDecl {{.*}} invalid gref '__global auto &__private' + __global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &__private' has incompatible initializer of type '__private int'}} __local int *ptr_l; - //CHECK: VarDecl {{.*}} invalid gptr '__global auto *' - __global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *' has incompatible initializer of type '__local int *'}} + //CHECK: VarDecl {{.*}} invalid gptr '__global auto *__private' + __global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *__private' has incompatible initializer of type '__local int *__private'}} }