[OpenCL] Add separate read_only and write_only pipe IR types

[stuart]

Summary

SPIR-V encodes the read_only and write_only access qualifiers of pipes, so separate LLVM IR types are required to target SPIR-V. Other backends may also find this useful.

These new types are opencl.pipe_ro_t and opencl.pipe_wo_t, which replace opencl.pipe_t.

This replaces __get_pipe_num_packets(...) and __get_pipe_max_packets(...) which took a read_only pipe with separate versions for read_only and write_only pipes, namely:

• __get_pipe_num_packets_ro(...)
• __get_pipe_num_packets_wo(...)
• __get_pipe_max_packets_ro(...)
• __get_pipe_max_packets_wo(...)

These separate versions exist to avoid needing a bitcast to one of the two qualified pipe types.

[AlexeySotkin]

It is not clear why we need two versions of get_pipe_num_packets and get_pipe_max_packets builtins. There is only one instruction per builtin in the SPIR-V spec. I think splitting the IR type is enough for translation to SPIR-V purposes.

[AlexeySotkin]

lib/CodeGen/CGOpenCLRuntime.h

65 ↗

(On Diff #143750)

I'm not sure that it is a good idea to make this function public, as its parameter supposed to be a reference to protected member.

[stuart]

Changed new getPipeType() method to have protected visibility.

Updated summary to explain the need for the extra builtin implementation functions.

[AlexeySotkin]

There should not be need for bitcast. Could give an example ? Thanks.

[stuart]

In D46015#1077401, @AlexeySotkin wrote:

It is not clear why we need two versions of get_pipe_num_packets and get_pipe_max_packets builtins. There is only one instruction per builtin in the SPIR-V spec. I think splitting the IR type is enough for translation to SPIR-V purposes.

This is so that when we emit the builtin expression, we can call a function that matches the access qualifier of the argument to the builtin, without the need for a bitcast of either the builtin's argument or the __get_pipe_max/num_packets() function itself.

lib/CodeGen/CGOpenCLRuntime.h

65 ↗

(On Diff #143750)

That's a good point. I have changed the function to be protected, to match the visibility of the data member.

[stuart]

In D46015#1078217, @AlexeySotkin wrote:

There should not be need for bitcast. Could give an example ? Thanks.

If I have a write_only pipe as the argument to get_pipe_max_packets(), and this uses a single __get_pipe_num_packets() function taking a read_only pipe, we will automatically get a bitcast:

%20 = call i32 bitcast (i32 (%opencl.pipe_ro_t*, i32, i32)* @__get_pipe_max_packets to i32 (%opencl.pipe_wo_t*, i32, i32)*)(%opencl.pipe_wo_t* %19, i32 4, i32 4)

[AlexeySotkin]

In D46015#1078235, @stuart wrote:

In D46015#1078217, @AlexeySotkin wrote:

There should not be need for bitcast. Could give an example ? Thanks.

If I have a write_only pipe as the argument to get_pipe_max_packets(), and this uses a single __get_pipe_num_packets() function taking a read_only pipe, we will automatically get a bitcast:

%20 = call i32 bitcast (i32 (%opencl.pipe_ro_t*, i32, i32)* @__get_pipe_max_packets to i32 (%opencl.pipe_wo_t*, i32, i32)*)(%opencl.pipe_wo_t* %19, i32 4, i32 4)

Sorry, but I don't quite understand what does get_pipe_max_packets(), uses __get_pipe_num_packets() mean. Could you clarify? Possibly OpenCL C source example could help.
Thanks

[stuart]

In D46015#1078260, @AlexeySotkin wrote:

In D46015#1078235, @stuart wrote:

In D46015#1078217, @AlexeySotkin wrote:

There should not be need for bitcast. Could give an example ? Thanks.

If I have a write_only pipe as the argument to get_pipe_max_packets(), and this uses a single __get_pipe_num_packets() function taking a read_only pipe, we will automatically get a bitcast:

%20 = call i32 bitcast (i32 (%opencl.pipe_ro_t*, i32, i32)* @__get_pipe_max_packets to i32 (%opencl.pipe_wo_t*, i32, i32)*)(%opencl.pipe_wo_t* %19, i32 4, i32 4)

Sorry, but I don't quite understand what does get_pipe_max_packets(), uses __get_pipe_num_packets() mean. Could you clarify? Possibly OpenCL C source example could help.

I mean that without these two separate versions, the call to __get_pipe_num_packets() that is emitted can include a bitcast.

For example:

void foo(read_only pipe int r, write_only pipe int w) {
  get_pipe_num_packets(w);
  get_pipe_num_packets(r);
}

get_pipe_num_packets(w) is seen first, causing i32 @__get_pipe_num_packets(%opencl.pipe_wo_t*, i32, i32) to be implicitly declared.

When the call to __get_pipe_num_packets() is emitted, this will be with an autogenerated bitcast from the type of the implicit declaration, i.e. i32 (%opencl.pipe_wo_t*, i32, i32)* to the type in the emitted expression, i.e. i32 (%opencl.pipe_ro_t*, i32, i32)*.

Here is the relevant section of IR:

%0 = load %opencl.pipe_wo_t*, %opencl.pipe_wo_t** %w.addr, align 8
%1 = call i32 @__get_pipe_num_packets(%opencl.pipe_wo_t* %0, i32 4, i32 4)
%2 = load %opencl.pipe_ro_t*, %opencl.pipe_ro_t** %r.addr, align 8
%3 = call i32 bitcast (i32 (%opencl.pipe_wo_t*, i32, i32)* @__get_pipe_num_packets to i32 (%opencl.pipe_ro_t*, i32, i32)*)(%opencl.pipe_ro_t* %2, i32 4, i32 4)

If we swap the two uses of get_pipe_num_packets() in the example above, then the type of the implicit declaration will be i32 (%opencl.pipe_ro_t*, i32, i32)* and bitcasts will instead be automatically generated when using get_pipe_num_packets() with a write_only pipe. It seems especially unfortunate that the type of the implicit declaration varies depending on the access qualifier of the first use.

[AlexeySotkin]

In D46015#1078317, @stuart wrote:

In D46015#1078260, @AlexeySotkin wrote:

In D46015#1078235, @stuart wrote:

In D46015#1078217, @AlexeySotkin wrote:

There should not be need for bitcast. Could give an example ? Thanks.

If I have a write_only pipe as the argument to get_pipe_max_packets(), and this uses a single __get_pipe_num_packets() function taking a read_only pipe, we will automatically get a bitcast:

%20 = call i32 bitcast (i32 (%opencl.pipe_ro_t*, i32, i32)* @__get_pipe_max_packets to i32 (%opencl.pipe_wo_t*, i32, i32)*)(%opencl.pipe_wo_t* %19, i32 4, i32 4)

Sorry, but I don't quite understand what does get_pipe_max_packets(), uses __get_pipe_num_packets() mean. Could you clarify? Possibly OpenCL C source example could help.

I mean that without these two separate versions, the call to __get_pipe_num_packets() that is emitted can include a bitcast.

For example:

void foo(read_only pipe int r, write_only pipe int w) {
  get_pipe_num_packets(w);
  get_pipe_num_packets(r);
}

get_pipe_num_packets(w) is seen first, causing i32 @__get_pipe_num_packets(%opencl.pipe_wo_t*, i32, i32) to be implicitly declared.

When the call to __get_pipe_num_packets() is emitted, this will be with an autogenerated bitcast from the type of the implicit declaration, i.e. i32 (%opencl.pipe_wo_t*, i32, i32)* to the type in the emitted expression, i.e. i32 (%opencl.pipe_ro_t*, i32, i32)*.

Here is the relevant section of IR:

%0 = load %opencl.pipe_wo_t*, %opencl.pipe_wo_t** %w.addr, align 8
%1 = call i32 @__get_pipe_num_packets(%opencl.pipe_wo_t* %0, i32 4, i32 4)
%2 = load %opencl.pipe_ro_t*, %opencl.pipe_ro_t** %r.addr, align 8
%3 = call i32 bitcast (i32 (%opencl.pipe_wo_t*, i32, i32)* @__get_pipe_num_packets to i32 (%opencl.pipe_ro_t*, i32, i32)*)(%opencl.pipe_ro_t* %2, i32 4, i32 4)

If we swap the two uses of get_pipe_num_packets() in the example above, then the type of the implicit declaration will be i32 (%opencl.pipe_ro_t*, i32, i32)* and bitcasts will instead be automatically generated when using get_pipe_num_packets() with a write_only pipe. It seems especially unfortunate that the type of the implicit declaration varies depending on the access qualifier of the first use.

Oh I see. LGTM then. Thanks.

cfe/trunk/lib/CodeGen/CGBuiltin.cpp

case Builtin::BIsub_group_commit_write_pipe: {
 
     return RValue::get(
         Builder.CreateCall(CGM.CreateRuntimeFunction(FTy, Name),
                            {Arg0, Arg1, PacketSize, PacketAlign}));
   }
   // OpenCL v2.0 s6.13.16.4 Built-in pipe query functions
   case Builtin::BIget_pipe_num_packets:
   case Builtin::BIget_pipe_max_packets: {
-    const char *Name;
+    const char *BaseName;
+    const PipeType *PipeTy = E->getArg(0)->getType()->getAs<PipeType>();
     if (BuiltinID == Builtin::BIget_pipe_num_packets)
-      Name = "__get_pipe_num_packets";
+      BaseName = "__get_pipe_num_packets";
     else
-      Name = "__get_pipe_max_packets";
+      BaseName = "__get_pipe_max_packets";
+    auto Name = std::string(BaseName) +
+                std::string(PipeTy->isReadOnly() ? "_ro" : "_wo");
 
     // Building the generic function prototype.
     Value *Arg0 = EmitScalarExpr(E->getArg(0));
     CGOpenCLRuntime OpenCLRT(CGM);
     Value *PacketSize = OpenCLRT.getPipeElemSize(E->getArg(0));
     Value *PacketAlign = OpenCLRT.getPipeElemAlign(E->getArg(0));
     llvm::Type *ArgTys[] = {Arg0->getType(), Int32Ty, Int32Ty};
     llvm::FunctionType *FTy = llvm::FunctionType::get(

cfe/trunk/lib/CodeGen/CGOpenCLRuntime.h

 namespace CodeGen {
 
 class CodeGenFunction;
 class CodeGenModule;
 
 class CGOpenCLRuntime {
 protected:
   CodeGenModule &CGM;
-  llvm::Type *PipeTy;
+  llvm::Type *PipeROTy;
+  llvm::Type *PipeWOTy;
   llvm::PointerType *SamplerTy;
 
   /// Structure for enqueued block information.
   struct EnqueuedBlockInfo {
     llvm::Function *InvokeFunc; /// Block invoke function.
     llvm::Function *Kernel;     /// Enqueued block kernel.
     llvm::Value *BlockArg;      /// The first argument to enqueued block kernel.
   };
   /// Maps block expression to block information.
   llvm::DenseMap<const Expr *, EnqueuedBlockInfo> EnqueuedBlockMap;
 
+  virtual llvm::Type *getPipeType(const PipeType *T, StringRef Name,
+                                  llvm::Type *&PipeTy);
+
 public:
-  CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM), PipeTy(nullptr),
-    SamplerTy(nullptr) {}
+  CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM),
+    PipeROTy(nullptr), PipeWOTy(nullptr), SamplerTy(nullptr) {}
   virtual ~CGOpenCLRuntime();
 
   /// Emit the IR required for a work-group-local variable declaration, and add
   /// an entry to CGF's LocalDeclMap for D.  The base class does this using
   /// CodeGenFunction::EmitStaticVarDecl to emit an internal global for D.
   virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
                                          const VarDecl &D);
 

cfe/trunk/lib/CodeGen/CGOpenCLRuntime.cpp

return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.queue_t"), AddrSpc);
   case BuiltinType::OCLReserveID:
     return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.reserve_id_t"), AddrSpc);
   }
 }
 
 llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) {
-  if (!PipeTy){
-    uint32_t PipeAddrSpc = CGM.getContext().getTargetAddressSpace(
-        CGM.getContext().getOpenCLTypeAddrSpace(T));
-    PipeTy = llvm::PointerType::get(llvm::StructType::create(
-      CGM.getLLVMContext(), "opencl.pipe_t"), PipeAddrSpc);
-  }
+  if (T->isReadOnly())
+    return getPipeType(T, "opencl.pipe_ro_t", PipeROTy);
+  else
+    return getPipeType(T, "opencl.pipe_wo_t", PipeWOTy);
 
+llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T, StringRef Name,
+                                         llvm::Type *&PipeTy) {
+  if (!PipeTy)
+    PipeTy = llvm::PointerType::get(llvm::StructType::create(
+      CGM.getLLVMContext(), Name),
+      CGM.getContext().getTargetAddressSpace(
+          CGM.getContext().getOpenCLTypeAddrSpace(T)));
   return PipeTy;
 }
 
 llvm::PointerType *CGOpenCLRuntime::getSamplerType(const Type *T) {
   if (!SamplerTy)
     SamplerTy = llvm::PointerType::get(llvm::StructType::create(
       CGM.getLLVMContext(), "opencl.sampler_t"),
       CGM.getContext().getTargetAddressSpace(

cfe/trunk/test/CodeGenOpenCL/opencl_types.cl

kernel void foo(image1d_t img) {
   fnc4smp(smp);
   // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
   fnc4smp(glb_smp);
   // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
 }
 
-kernel void foo_pipe(read_only pipe int p) {}
-// CHECK-SPIR: @foo_pipe(%opencl.pipe_t addrspace(1)* %p)
-// CHECK_AMDGCN: @foo_pipe(%opencl.pipe_t addrspace(1)* %p)
+kernel void foo_ro_pipe(read_only pipe int p) {}
+// CHECK-SPIR: @foo_ro_pipe(%opencl.pipe_ro_t addrspace(1)* %p)
+// CHECK_AMDGCN: @foo_ro_pipe(%opencl.pipe_ro_t addrspace(1)* %p)
+
+kernel void foo_wo_pipe(write_only pipe int p) {}
+// CHECK-SPIR: @foo_wo_pipe(%opencl.pipe_wo_t addrspace(1)* %p)
+// CHECK_AMDGCN: @foo_wo_pipe(%opencl.pipe_wo_t addrspace(1)* %p)
 
 void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
 // CHECK-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}
 // CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(%opencl.image1d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}})

cfe/trunk/test/CodeGenOpenCL/pipe_builtin.cl

 // RUN: %clang_cc1 -emit-llvm -cl-ext=+cl_khr_subgroups -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
 
-// CHECK: %opencl.pipe_t = type opaque
-// CHECK: %opencl.reserve_id_t = type opaque
+// CHECK-DAG: %opencl.pipe_ro_t = type opaque
+// CHECK-DAG: %opencl.pipe_wo_t = type opaque
+// CHECK-DAG: %opencl.reserve_id_t = type opaque
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
 void test1(read_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
   read_pipe(p, ptr);
-  // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = reserve_read_pipe(p, 2);
-  // CHECK: call i32 @__read_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__read_pipe_4(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
   read_pipe(p, rid, 2, ptr);
-  // CHECK: call void @__commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   commit_read_pipe(p, rid);
 }
 
 void test2(write_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__write_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__write_pipe_2(%opencl.pipe_wo_t* %{{.*}}, i8* %{{.*}}, i32 4, i32 4)
   write_pipe(p, ptr);
-  // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = reserve_write_pipe(p, 2);
-  // CHECK: call i32 @__write_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__write_pipe_4(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}}, i32 4, i32 4)
   write_pipe(p, rid, 2, ptr);
-  // CHECK: call void @__commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   commit_write_pipe(p, rid);
 }
 
 void test3(read_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = work_group_reserve_read_pipe(p, 2);
-  // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   work_group_commit_read_pipe(p, rid);
 }
 
 void test4(write_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = work_group_reserve_write_pipe(p, 2);
-  // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   work_group_commit_write_pipe(p, rid);
 }
 
 void test5(read_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_ro_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = sub_group_reserve_read_pipe(p, 2);
-  // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_ro_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   sub_group_commit_read_pipe(p, rid);
 }
 
 void test6(write_only pipe int p, global int *ptr) {
-  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
+  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_wo_t* %{{.*}}, i32 {{.*}}, i32 4, i32 4)
   reserve_id_t rid = sub_group_reserve_write_pipe(p, 2);
-  // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_wo_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 4, i32 4)
   sub_group_commit_write_pipe(p, rid);
 }
 
-void test7(write_only pipe int p, global int *ptr) {
-  // CHECK: call i32 @__get_pipe_num_packets(%opencl.pipe_t* %{{.*}}, i32 4, i32 4)
+void test7(read_only pipe int p, global int *ptr) {
+  // CHECK: call i32 @__get_pipe_num_packets_ro(%opencl.pipe_ro_t* %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_num_packets(p);
-  // CHECK: call i32 @__get_pipe_max_packets(%opencl.pipe_t* %{{.*}}, i32 4, i32 4)
+  // CHECK: call i32 @__get_pipe_max_packets_ro(%opencl.pipe_ro_t* %{{.*}}, i32 4, i32 4)
   *ptr = get_pipe_max_packets(p);
 }
 
-void test8(read_only pipe int r, write_only pipe int w, global int *ptr) {
+void test8(write_only pipe int p, global int *ptr) {
+  // CHECK: call i32 @__get_pipe_num_packets_wo(%opencl.pipe_wo_t* %{{.*}}, i32 4, i32 4)
+  *ptr = get_pipe_num_packets(p);
+  // CHECK: call i32 @__get_pipe_max_packets_wo(%opencl.pipe_wo_t* %{{.*}}, i32 4, i32 4)
+  *ptr = get_pipe_max_packets(p);
+}
+
+void test9(read_only pipe int r, write_only pipe int w, global int *ptr) {
   // verify that return type is correctly casted to i1 value
   // CHECK: %[[R:[0-9]+]] = call i32 @__read_pipe_2
   // CHECK: icmp ne i32 %[[R]], 0
   if (read_pipe(r, ptr)) *ptr = -1;
   // CHECK: %[[W:[0-9]+]] = call i32 @__write_pipe_2
   // CHECK: icmp ne i32 %[[W]], 0
   if (write_pipe(w, ptr)) *ptr = -1;
-  // CHECK: %[[N:[0-9]+]] = call i32 @__get_pipe_num_packets
-  // CHECK: icmp ne i32 %[[N]], 0
+  // CHECK: %[[NR:[0-9]+]] = call i32 @__get_pipe_num_packets_ro
+  // CHECK: icmp ne i32 %[[NR]], 0
   if (get_pipe_num_packets(r)) *ptr = -1;
-  // CHECK: %[[M:[0-9]+]] = call i32 @__get_pipe_max_packets
-  // CHECK: icmp ne i32 %[[M]], 0
+  // CHECK: %[[NW:[0-9]+]] = call i32 @__get_pipe_num_packets_wo
+  // CHECK: icmp ne i32 %[[NW]], 0
+  if (get_pipe_num_packets(w)) *ptr = -1;
+  // CHECK: %[[MR:[0-9]+]] = call i32 @__get_pipe_max_packets_ro
+  // CHECK: icmp ne i32 %[[MR]], 0
+  if (get_pipe_max_packets(r)) *ptr = -1;
+  // CHECK: %[[MW:[0-9]+]] = call i32 @__get_pipe_max_packets_wo
+  // CHECK: icmp ne i32 %[[MW]], 0
   if (get_pipe_max_packets(w)) *ptr = -1;
 }

cfe/trunk/test/CodeGenOpenCL/pipe_types.cl

 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
 
-// CHECK: %opencl.pipe_t = type opaque
+// CHECK: %opencl.pipe_ro_t = type opaque
+// CHECK: %opencl.pipe_wo_t = type opaque
 typedef unsigned char __attribute__((ext_vector_type(3))) uchar3;
 typedef int __attribute__((ext_vector_type(4))) int4;
 
 void test1(read_only pipe int p) {
-// CHECK: define void @test1(%opencl.pipe_t* %p)
+// CHECK: define void @test1(%opencl.pipe_ro_t* %p)
   reserve_id_t rid;
 // CHECK: %rid = alloca %opencl.reserve_id_t
 }
 
 void test2(write_only pipe float p) {
-// CHECK: define void @test2(%opencl.pipe_t* %p)
+// CHECK: define void @test2(%opencl.pipe_wo_t* %p)
 }
 
 void test3(read_only pipe const int p) {
-// CHECK: define void @test3(%opencl.pipe_t* %p)
+// CHECK: define void @test3(%opencl.pipe_ro_t* %p)
 }
 
 void test4(read_only pipe uchar3 p) {
-// CHECK: define void @test4(%opencl.pipe_t* %p)
+// CHECK: define void @test4(%opencl.pipe_ro_t* %p)
 }
 
 void test5(read_only pipe int4 p) {
-// CHECK: define void @test5(%opencl.pipe_t* %p)
+// CHECK: define void @test5(%opencl.pipe_ro_t* %p)
 }
 
 typedef read_only pipe int MyPipe;
 kernel void test6(MyPipe p) {
-// CHECK: define spir_kernel void @test6(%opencl.pipe_t* %p)
+// CHECK: define spir_kernel void @test6(%opencl.pipe_ro_t* %p)
 }
 
 struct Person {
   const char *Name;
   bool isFemale;
   int ID;
 };
 
 void test_reserved_read_pipe(global struct Person *SDst,
                              read_only pipe struct Person SPipe) {
 // CHECK: define void @test_reserved_read_pipe
   read_pipe (SPipe, SDst);
-  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
   read_pipe (SPipe, SDst);
-  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_ro_t* %{{.*}}, i8* %{{.*}}, i32 16, i32 8)
 }

cfe/trunk/test/Index/pipe-size.cl

 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN
 __kernel void testPipe( pipe int test )
 {
     int s = sizeof(test);
-    // X86: store %opencl.pipe_t* %test, %opencl.pipe_t** %test.addr, align 8
+    // X86: store %opencl.pipe_ro_t* %test, %opencl.pipe_ro_t** %test.addr, align 8
     // X86: store i32 8, i32* %s, align 4
-    // SPIR: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 4
+    // SPIR: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 4
     // SPIR: store i32 4, i32* %s, align 4
-    // SPIR64: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 8
+    // SPIR64: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 8
     // SPIR64: store i32 8, i32* %s, align 4
-    // AMDGCN: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 8
+    // AMDGCN: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)* addrspace(5)* %test.addr, align 8
     // AMDGCN: store i32 8, i32 addrspace(5)* %s, align 4
 }