echuraev created this revision. echuraev added a reviewer: Anastasia. echuraev added subscribers: cfe-commits, yaxunl, bader.
https://reviews.llvm.org/D27671 Files: include/clang/Basic/DiagnosticSemaKinds.td lib/Sema/SemaDecl.cpp test/SemaOpenCL/invalid-kernel.cl Index: test/SemaOpenCL/invalid-kernel.cl =================================================================== --- test/SemaOpenCL/invalid-kernel.cl +++ test/SemaOpenCL/invalid-kernel.cl @@ -1,10 +1,11 @@ // RUN: %clang_cc1 -verify %s +// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s -kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}} +kernel void no_ptrptr(global int * global *i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}} -__kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}} +__kernel void no_privateptr(__private int *i) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} -__kernel void no_privatearray(__private int i[]) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}} +__kernel void no_privatearray(__private int i[]) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} kernel int bar() { // expected-error {{kernel must have void return type}} return 6; @@ -29,3 +30,6 @@ int* constant x(int* x) { // expected-error {{return value cannot be qualified with address space}} return x + 1; } + +__kernel void testKernel(int *ptr) { // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} +} Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -7586,7 +7586,7 @@ ValidKernelParam, PtrPtrKernelParam, PtrKernelParam, - PrivatePtrKernelParam, + InvalidAddrSpacePtrKernelParam, InvalidKernelParam, RecordKernelParam }; @@ -7596,8 +7596,10 @@ QualType PointeeType = PT->getPointeeType(); if (PointeeType->isPointerType()) return PtrPtrKernelParam; - return PointeeType.getAddressSpace() == 0 ? PrivatePtrKernelParam - : PtrKernelParam; + if (PointeeType.getAddressSpace() == LangAS::opencl_generic || + PointeeType.getAddressSpace() == 0) + return InvalidAddrSpacePtrKernelParam; + return PtrKernelParam; } // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can @@ -7645,11 +7647,12 @@ D.setInvalidType(); return; - case PrivatePtrKernelParam: - // OpenCL v1.2 s6.9.a: - // A kernel function argument cannot be declared as a - // pointer to the private address space. - S.Diag(Param->getLocation(), diag::err_opencl_private_ptr_kernel_param); + case InvalidAddrSpacePtrKernelParam: + // OpenCL v1.0 s6.5: + // __kernel function arguments declared to be a pointer of a type can point + // to one of the following address spaces only : __global, __local or + // __constant. + S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space); D.setInvalidType(); return; @@ -7736,7 +7739,7 @@ // do not allow OpenCL objects to be passed as elements of the struct or // union. if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam || - ParamType == PrivatePtrKernelParam) { + ParamType == InvalidAddrSpacePtrKernelParam) { S.Diag(Param->getLocation(), diag::err_record_with_pointers_kernel_param) << PT->isUnionType() Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -8076,8 +8076,9 @@ "kernel functions cannot be declared static">; def err_opencl_ptrptr_kernel_param : Error< "kernel parameter cannot be declared as a pointer to a pointer">; -def err_opencl_private_ptr_kernel_param : Error< - "kernel parameter cannot be declared as a pointer to the __private address space">; +def err_kernel_arg_address_space : Error< + "pointer arguments to kernel functions must reside in '__global', " + "'__constant' or '__local' address space">; def err_opencl_function_variable : Error< "%select{non-kernel function|function scope}0 variable cannot be declared in %1 address space">; def err_static_function_scope : Error<
Index: test/SemaOpenCL/invalid-kernel.cl =================================================================== --- test/SemaOpenCL/invalid-kernel.cl +++ test/SemaOpenCL/invalid-kernel.cl @@ -1,10 +1,11 @@ // RUN: %clang_cc1 -verify %s +// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s -kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}} +kernel void no_ptrptr(global int * global *i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}} -__kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}} +__kernel void no_privateptr(__private int *i) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} -__kernel void no_privatearray(__private int i[]) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}} +__kernel void no_privatearray(__private int i[]) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} kernel int bar() { // expected-error {{kernel must have void return type}} return 6; @@ -29,3 +30,6 @@ int* constant x(int* x) { // expected-error {{return value cannot be qualified with address space}} return x + 1; } + +__kernel void testKernel(int *ptr) { // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}} +} Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -7586,7 +7586,7 @@ ValidKernelParam, PtrPtrKernelParam, PtrKernelParam, - PrivatePtrKernelParam, + InvalidAddrSpacePtrKernelParam, InvalidKernelParam, RecordKernelParam }; @@ -7596,8 +7596,10 @@ QualType PointeeType = PT->getPointeeType(); if (PointeeType->isPointerType()) return PtrPtrKernelParam; - return PointeeType.getAddressSpace() == 0 ? PrivatePtrKernelParam - : PtrKernelParam; + if (PointeeType.getAddressSpace() == LangAS::opencl_generic || + PointeeType.getAddressSpace() == 0) + return InvalidAddrSpacePtrKernelParam; + return PtrKernelParam; } // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can @@ -7645,11 +7647,12 @@ D.setInvalidType(); return; - case PrivatePtrKernelParam: - // OpenCL v1.2 s6.9.a: - // A kernel function argument cannot be declared as a - // pointer to the private address space. - S.Diag(Param->getLocation(), diag::err_opencl_private_ptr_kernel_param); + case InvalidAddrSpacePtrKernelParam: + // OpenCL v1.0 s6.5: + // __kernel function arguments declared to be a pointer of a type can point + // to one of the following address spaces only : __global, __local or + // __constant. + S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space); D.setInvalidType(); return; @@ -7736,7 +7739,7 @@ // do not allow OpenCL objects to be passed as elements of the struct or // union. if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam || - ParamType == PrivatePtrKernelParam) { + ParamType == InvalidAddrSpacePtrKernelParam) { S.Diag(Param->getLocation(), diag::err_record_with_pointers_kernel_param) << PT->isUnionType() Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -8076,8 +8076,9 @@ "kernel functions cannot be declared static">; def err_opencl_ptrptr_kernel_param : Error< "kernel parameter cannot be declared as a pointer to a pointer">; -def err_opencl_private_ptr_kernel_param : Error< - "kernel parameter cannot be declared as a pointer to the __private address space">; +def err_kernel_arg_address_space : Error< + "pointer arguments to kernel functions must reside in '__global', " + "'__constant' or '__local' address space">; def err_opencl_function_variable : Error< "%select{non-kernel function|function scope}0 variable cannot be declared in %1 address space">; def err_static_function_scope : Error<
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits