[OpenCL] Add separate read_only and write_only pipe IR types

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.

Patch by Stuart Brady.

Differential Revision: https://reviews.llvm.org/D46015

llvm-svn: 331026

diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 6a2f2b0..3987db9 100644 (file)
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -3051,11 +3051,14 @@ RValue CodeGenFunction::EmitBuiltinExpr(const FunctionDecl *FD,
   // 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));

diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
index 15a2221..1da19a9 100644 (file)
--- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
@@ -66,13 +66,19 @@ llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
 }
 
 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;
 }
 

diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.h b/clang/lib/CodeGen/CGOpenCLRuntime.h
index 9378b62..2df1acd 100644 (file)
--- a/clang/lib/CodeGen/CGOpenCLRuntime.h
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.h
@@ -35,7 +35,8 @@ class CodeGenModule;
 class CGOpenCLRuntime {
 protected:
   CodeGenModule &CGM;
-  llvm::Type *PipeTy;
+  llvm::Type *PipeROTy;
+  llvm::Type *PipeWOTy;
   llvm::PointerType *SamplerTy;
 
   /// Structure for enqueued block information.
@@ -47,9 +48,12 @@ protected:
   /// 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

diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl
index ee3c8d7..1f1294b 100644 (file)
--- a/clang/test/CodeGenOpenCL/opencl_types.cl
+++ b/clang/test/CodeGenOpenCL/opencl_types.cl
@@ -63,9 +63,13 @@ kernel void foo(image1d_t img) {
   // 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"}}

diff --git a/clang/test/CodeGenOpenCL/pipe_builtin.cl b/clang/test/CodeGenOpenCL/pipe_builtin.cl
index 4b4b4ef..d912fce 100644 (file)
--- a/clang/test/CodeGenOpenCL/pipe_builtin.cl
+++ b/clang/test/CodeGenOpenCL/pipe_builtin.cl
@@ -1,68 +1,76 @@
 // 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
@@ -70,10 +78,16 @@ void test8(read_only pipe int r, write_only pipe int w, global int *ptr) {
   // 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;
 }

diff --git a/clang/test/CodeGenOpenCL/pipe_types.cl b/clang/test/CodeGenOpenCL/pipe_types.cl
index 7c11f74..ba064c6 100644 (file)
--- a/clang/test/CodeGenOpenCL/pipe_types.cl
+++ b/clang/test/CodeGenOpenCL/pipe_types.cl
@@ -1,34 +1,35 @@
 // 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 {
@@ -41,7 +42,7 @@ 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)
 }

diff --git a/clang/test/Index/pipe-size.cl b/clang/test/Index/pipe-size.cl
index 78a8f43..94a1255 100644 (file)
--- a/clang/test/Index/pipe-size.cl
+++ b/clang/test/Index/pipe-size.cl
@@ -5,12 +5,12 @@
 __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
 }