dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[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


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331026 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Sven van Haastregt 2018-04-27 10:37:04 +00:00
parent a2f75a9805
commit 2bf83a47cc
7 changed files with 86 additions and 54 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
lib/CodeGen/CGBuiltin.cpp
View File

@ -3051,11 +3051,14 @@ RValue CodeGenFunction::EmitBuiltinExpr(const FunctionDecl *FD,
// OpenCL v2.0 s6.13.16.4 Built-in pipe query functions // OpenCL v2.0 s6.13.16.4 Built-in pipe query functions
case Builtin::BIget_pipe_num_packets: case Builtin::BIget_pipe_num_packets:
case Builtin::BIget_pipe_max_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) if (BuiltinID == Builtin::BIget_pipe_num_packets)
Name = "__get_pipe_num_packets"; BaseName = "__get_pipe_num_packets";
else 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. // Building the generic function prototype.
Value *Arg0 = EmitScalarExpr(E->getArg(0)); Value *Arg0 = EmitScalarExpr(E->getArg(0));

18
lib/CodeGen/CGOpenCLRuntime.cpp
View File

@ -66,13 +66,19 @@ llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
} }
llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) { llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) {
if (!PipeTy){ if (T->isReadOnly())
uint32_t PipeAddrSpc = CGM.getContext().getTargetAddressSpace( return getPipeType(T, "opencl.pipe_ro_t", PipeROTy);
CGM.getContext().getOpenCLTypeAddrSpace(T)); else
PipeTy = llvm::PointerType::get(llvm::StructType::create( return getPipeType(T, "opencl.pipe_wo_t", PipeWOTy);
CGM.getLLVMContext(), "opencl.pipe_t"), PipeAddrSpc); }
}
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; return PipeTy;
} }

10
lib/CodeGen/CGOpenCLRuntime.h
View File

@ -35,7 +35,8 @@ class CodeGenModule;
class CGOpenCLRuntime { class CGOpenCLRuntime {
protected: protected:
CodeGenModule &CGM; CodeGenModule &CGM;
llvm::Type *PipeTy; llvm::Type *PipeROTy;
llvm::Type *PipeWOTy;
llvm::PointerType *SamplerTy; llvm::PointerType *SamplerTy;
/// Structure for enqueued block information. /// Structure for enqueued block information.
@ -47,9 +48,12 @@ protected:
/// Maps block expression to block information. /// Maps block expression to block information.
llvm::DenseMap<const Expr *, EnqueuedBlockInfo> EnqueuedBlockMap; llvm::DenseMap<const Expr *, EnqueuedBlockInfo> EnqueuedBlockMap;
virtual llvm::Type *getPipeType(const PipeType *T, StringRef Name,
llvm::Type *&PipeTy);
public: public:
CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM), PipeTy(nullptr), CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM),
SamplerTy(nullptr) {} PipeROTy(nullptr), PipeWOTy(nullptr), SamplerTy(nullptr) {}
virtual ~CGOpenCLRuntime(); virtual ~CGOpenCLRuntime();
/// Emit the IR required for a work-group-local variable declaration, and add /// Emit the IR required for a work-group-local variable declaration, and add

10
test/CodeGenOpenCL/opencl_types.cl
View File

@ -63,9 +63,13 @@ kernel void foo(image1d_t img) {
// CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)* // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
} }
kernel void foo_pipe(read_only pipe int p) {} kernel void foo_ro_pipe(read_only pipe int p) {}
// CHECK-SPIR: @foo_pipe(%opencl.pipe_t addrspace(1)* %p) // CHECK-SPIR: @foo_ro_pipe(%opencl.pipe_ro_t addrspace(1)* %p)
// CHECK_AMDGCN: @foo_pipe(%opencl.pipe_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) {} 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-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}

66
test/CodeGenOpenCL/pipe_builtin.cl
View File

@ -1,68 +1,76 @@
// RUN: %clang_cc1 -emit-llvm -cl-ext=+cl_khr_subgroups -O0 -cl-std=CL2.0 -o - %s | FileCheck %s // 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-DAG: %opencl.pipe_ro_t = type opaque
// CHECK: %opencl.reserve_id_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 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
void test1(read_only pipe int p, global int *ptr) { 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); 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); 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); 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); commit_read_pipe(p, rid);
} }
void test2(write_only pipe int p, global int *ptr) { 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); 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); 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); 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); commit_write_pipe(p, rid);
} }
void test3(read_only pipe int p, global int *ptr) { 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); 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); work_group_commit_read_pipe(p, rid);
} }
void test4(write_only pipe int p, global int *ptr) { 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); 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); work_group_commit_write_pipe(p, rid);
} }
void test5(read_only pipe int p, global int *ptr) { 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); 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); sub_group_commit_read_pipe(p, rid);
} }
void test6(write_only pipe int p, global int *ptr) { 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); 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); sub_group_commit_write_pipe(p, rid);
} }
void test7(write_only pipe int p, global int *ptr) { void test7(read_only pipe int p, global int *ptr) {
// CHECK: call i32 @__get_pipe_num_packets(%opencl.pipe_t* %{{.*}}, i32 4, i32 4) // CHECK: call i32 @__get_pipe_num_packets_ro(%opencl.pipe_ro_t* %{{.*}}, i32 4, i32 4)
*ptr = get_pipe_num_packets(p); *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); *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 // verify that return type is correctly casted to i1 value
// CHECK: %[[R:[0-9]+]] = call i32 @__read_pipe_2 // CHECK: %[[R:[0-9]+]] = call i32 @__read_pipe_2
// CHECK: icmp ne i32 %[[R]], 0 // 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: %[[W:[0-9]+]] = call i32 @__write_pipe_2
// CHECK: icmp ne i32 %[[W]], 0 // CHECK: icmp ne i32 %[[W]], 0
if (write_pipe(w, ptr)) *ptr = -1; if (write_pipe(w, ptr)) *ptr = -1;
// CHECK: %[[N:[0-9]+]] = call i32 @__get_pipe_num_packets // CHECK: %[[NR:[0-9]+]] = call i32 @__get_pipe_num_packets_ro
// CHECK: icmp ne i32 %[[N]], 0 // CHECK: icmp ne i32 %[[NR]], 0
if (get_pipe_num_packets(r)) *ptr = -1; if (get_pipe_num_packets(r)) *ptr = -1;
// CHECK: %[[M:[0-9]+]] = call i32 @__get_pipe_max_packets // CHECK: %[[NW:[0-9]+]] = call i32 @__get_pipe_num_packets_wo
// CHECK: icmp ne i32 %[[M]], 0 // 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; if (get_pipe_max_packets(w)) *ptr = -1;
} }

19
test/CodeGenOpenCL/pipe_types.cl
View File

@ -1,34 +1,35 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s // 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 unsigned char __attribute__((ext_vector_type(3))) uchar3;
typedef int __attribute__((ext_vector_type(4))) int4; typedef int __attribute__((ext_vector_type(4))) int4;
void test1(read_only pipe int p) { 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; reserve_id_t rid;
// CHECK: %rid = alloca %opencl.reserve_id_t // CHECK: %rid = alloca %opencl.reserve_id_t
} }
void test2(write_only pipe float p) { 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) { 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) { 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) { 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; typedef read_only pipe int MyPipe;
kernel void test6(MyPipe p) { 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 { struct Person {
@ -41,7 +42,7 @@ void test_reserved_read_pipe(global struct Person *SDst,
read_only pipe struct Person SPipe) { read_only pipe struct Person SPipe) {
// CHECK: define void @test_reserved_read_pipe // CHECK: define void @test_reserved_read_pipe
read_pipe (SPipe, SDst); 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); 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)
} }

8
test/Index/pipe-size.cl
View File

@ -5,12 +5,12 @@
__kernel void testPipe( pipe int test ) __kernel void testPipe( pipe int test )
{ {
int s = sizeof(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 // 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 // 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 // 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 // AMDGCN: store i32 8, i32 addrspace(5)* %s, align 4
} }