diff options
| author | Sven van Haastregt <sven.vanhaastregt@arm.com> | 2018-04-27 10:37:04 +0000 |
|---|---|---|
| committer | Sven van Haastregt <sven.vanhaastregt@arm.com> | 2018-04-27 10:37:04 +0000 |
| commit | 2bf83a47cc2912d161c8d2240b7816cc61ea8a10 (patch) | |
| tree | 80e984151f66eaeff8f50fb46d6c2444268e766d /test/Index | |
| parent | a2f75a9805a64c0674b3d9cffdf22c72fbb71ddd (diff) | |
[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
Diffstat (limited to 'test/Index')
| -rw-r--r-- | test/Index/pipe-size.cl | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/test/Index/pipe-size.cl b/test/Index/pipe-size.cl index 78a8f43810..94a1255f0a 100644 --- a/test/Index/pipe-size.cl +++ b/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 } |