Clang error: invalid pipe access modifier (expecting A) (err_opencl_builtin_pipe_invalid_access_modifier)

Text	error: invalid pipe access modifier (expecting A)
Type	Error
Category	Semantic Issue
Internal Id	err_opencl_builtin_pipe_invalid_access_modifier
Internal Message	invalid pipe access modifier (expecting %0)
Regular Expression	(?:error|fatal error)\: invalid pipe access modifier \(expecting (.*?)\)
First Commit	2016-01-26 3a9952c9e7d0 [OpenCL] Pipe builtin functions

Description

The error is issued by the Clang compiler when attempting to use OpenCL pipe built-in functions with an incorrect access modifier. OpenCL 2.0 introduces pipes as a mechanism for FIFO communication between different device kernels. These pipes are declared with access qualifiers read_only and write_only to specify how the kernel intends to interact with the pipe. The error occurs if a pipe operation is called with a pipe object whose access qualifier does not match the operation's expected access mode. For example, using read_pipe or any other reading operation on a pipe declared as write_only triggers this error because the operation expects the pipe to be accessible for reading, not just writing. Conversely, attempting to write to a read_only pipe using operations like write_pipe will also result in this error. The Clang compiler checks the access qualifier of the pipe argument against the expected modifier of the pipe built-in function being called to ensure proper usage and data flow in OpenCL applications.

AI Generated

Example

In the following example, an OpenCL kernel function is defined, which is intended to perform operations on a pipe object. The pipe object is declared with a __write_only access qualifier, indicating that operations on this pipe should only write data to it. The kernel then attempts to perform a read operation on this write-only pipe by calling a built-in function designed for reading from pipes. This mismatch between the pipe's access qualifier and the operation attempting to access it in a manner contrary to its declaration triggers the Clang error message. The error output explicitly states the issue by pointing out the invalid access modifier and indicating the correct expected access modifier for the operation being performed.

AI Generated

Flags	-xcl -cl-std=CL2.0 -Xclang -finclude-default-header	[Try out in Compiler Explorer]
Source	__kernel void test(__write_only pipe int p) { // Incorrect pipe access modifier int v; sub_group_commit_read_pipe(p, &v); }
Compiler Output	<source>:3:30: error: invalid pipe access modifier (expecting read_only)

Clang Internals (17.0.6)

Git Commit Message

[OpenCL] Pipe builtin functions

Summary:
Support for the pipe built-in functions for OpenCL 2.0.
The pipe builtin functions may have infinite kinds of element types, one approach
would be to just generate calls that would always use generic types such as void*.
This patch is based on bader's opencl support patch on SPIR-V branch.

Reviewers: Anastasia, pekka.jaaskelainen

Subscribers: keryell, bader, cfe-commits

Differential Revision: http://reviews.llvm.org/D15914

llvm-svn: 258773

Used in Clang Sources

This section lists all occurrences of the diagnostic within the Clang's codebase. For each occurrence, an auto-extracted snipped from the source code is listed including key elements like control structures, functions, or classes. It should illustrate the conditions under which the diagnostic is activated.

clang/lib/Sema/SemaChecking.cpp (line 1677)

/// Returns true if pipe element type is different from the pointer.
static bool checkOpenCLPipeArg(Sema &S, CallExpr *Call) {
  // ...
  case Builtin::BIread_pipe:
  case Builtin::BIreserve_read_pipe:
  case Builtin::BIcommit_read_pipe:
  case Builtin::BIwork_group_reserve_read_pipe:
  case Builtin::BIsub_group_reserve_read_pipe:
  case Builtin::BIwork_group_commit_read_pipe:
  case Builtin::BIsub_group_commit_read_pipe:
    if (!(!AccessQual || AccessQual->isReadOnly())) {
      S.Diag(Arg0->getBeginLoc(), diag::err_opencl_builtin_pipe_invalid_access_modifier) << "read_only" << Arg0->getSourceRange();

clang/lib/Sema/SemaChecking.cpp (line 1691)

/// Returns true if pipe element type is different from the pointer.
static bool checkOpenCLPipeArg(Sema &S, CallExpr *Call) {
  // ...
  case Builtin::BIwrite_pipe:
  case Builtin::BIreserve_write_pipe:
  case Builtin::BIcommit_write_pipe:
  case Builtin::BIwork_group_reserve_write_pipe:
  case Builtin::BIsub_group_reserve_write_pipe:
  case Builtin::BIwork_group_commit_write_pipe:
  case Builtin::BIsub_group_commit_write_pipe:
    if (!(AccessQual && AccessQual->isWriteOnly())) {
      S.Diag(Arg0->getBeginLoc(), diag::err_opencl_builtin_pipe_invalid_access_modifier) << "write_only" << Arg0->getSourceRange();

Triggered in Clang Tests

This section lists all internal Clang test cases that trigger the diagnostic.

clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:18:14: error: invalid pipe access modifier (expecting write_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:19:14: error: invalid pipe access modifier (expecting write_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:25:32: error: invalid pipe access modifier (expecting write_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:31:31: error: invalid pipe access modifier (expecting write_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:47:13: error: invalid pipe access modifier (expecting read_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:48:13: error: invalid pipe access modifier (expecting read_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:54:31: error: invalid pipe access modifier (expecting read_only)

• clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl:60:30: error: invalid pipe access modifier (expecting read_only)