Index: lib/CodeGen/CGBuiltin.cpp =================================================================== --- lib/CodeGen/CGBuiltin.cpp +++ lib/CodeGen/CGBuiltin.cpp @@ -3051,11 +3051,14 @@ // 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(); 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)); Index: lib/CodeGen/CGOpenCLRuntime.h =================================================================== --- lib/CodeGen/CGOpenCLRuntime.h +++ lib/CodeGen/CGOpenCLRuntime.h @@ -35,7 +35,8 @@ 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 @@ /// Maps block expression to block information. llvm::DenseMap 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 Index: lib/CodeGen/CGOpenCLRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenCLRuntime.cpp +++ lib/CodeGen/CGOpenCLRuntime.cpp @@ -66,13 +66,19 @@ } 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; } Index: test/CodeGenOpenCL/opencl_types.cl =================================================================== --- test/CodeGenOpenCL/opencl_types.cl +++ test/CodeGenOpenCL/opencl_types.cl @@ -63,9 +63,13 @@ // 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"}} Index: test/CodeGenOpenCL/pipe_builtin.cl =================================================================== --- test/CodeGenOpenCL/pipe_builtin.cl +++ 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 @@ // 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; } Index: test/CodeGenOpenCL/pipe_types.cl =================================================================== --- test/CodeGenOpenCL/pipe_types.cl +++ 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 @@ 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) } Index: test/Index/pipe-size.cl =================================================================== --- test/Index/pipe-size.cl +++ 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 }