1 ; The following SPIR 2.0 was obtained via SPIR-V generator/Clang: 2 ; bash$ clang -cc1 -x cl -cl-std=CL2.0 -triple spir64-unknonw-unknown -emit-llvm -include opencl-20.h -Dcl_khr_subgroups pipe_builtins.cl -o pipe_builtins.ll 3 4 ;; Regression test: 5 ;; Pipe built-ins are mangled accordingly to SPIR2.0/C++ ABI. 6 7 ; #pragma OPENCL EXTENSION cl_khr_subgroups : enable 8 ; 9 ; __kernel void test_pipe_convenience_write_uint(__global uint *src, __write_only pipe uint out_pipe) 10 ; { 11 ; int gid = get_global_id(0); 12 ; write_pipe(out_pipe, &src[gid]); 13 ; } 14 ; 15 ; __kernel void test_pipe_convenience_read_uint(__read_only pipe uint in_pipe, __global uint *dst) 16 ; { 17 ; int gid = get_global_id(0); 18 ; read_pipe(in_pipe, &dst[gid]); 19 ; } 20 ; 21 ; __kernel void test_pipe_write(__global int *src, __write_only pipe int out_pipe) 22 ; { 23 ; int gid = get_global_id(0); 24 ; reserve_id_t res_id; 25 ; res_id = reserve_write_pipe(out_pipe, 1); 26 ; if(is_valid_reserve_id(res_id)) 27 ; { 28 ; write_pipe(out_pipe, res_id, 0, &src[gid]); 29 ; commit_write_pipe(out_pipe, res_id); 30 ; } 31 ; } 32 ; 33 ; __kernel void test_pipe_query_functions(__write_only pipe int out_pipe, __global int *num_packets, __global int *max_packets) 34 ; { 35 ; *max_packets = get_pipe_max_packets(out_pipe); 36 ; *num_packets = get_pipe_num_packets(out_pipe); 37 ; } 38 ; 39 ; __kernel void test_pipe_read(__read_only pipe int in_pipe, __global int *dst) 40 ; { 41 ; int gid = get_global_id(0); 42 ; reserve_id_t res_id; 43 ; res_id = reserve_read_pipe(in_pipe, 1); 44 ; if(is_valid_reserve_id(res_id)) 45 ; { 46 ; read_pipe(in_pipe, res_id, 0, &dst[gid]); 47 ; commit_read_pipe(in_pipe, res_id); 48 ; } 49 ; } 50 ; 51 ; __kernel void test_pipe_workgroup_write_char(__global char *src, __write_only pipe char out_pipe) 52 ; { 53 ; int gid = get_global_id(0); 54 ; __local reserve_id_t res_id; 55 ; 56 ; res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0)); 57 ; if(is_valid_reserve_id(res_id)) 58 ; { 59 ; write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]); 60 ; work_group_commit_write_pipe(out_pipe, res_id); 61 ; } 62 ; } 63 ; 64 ; __kernel void test_pipe_workgroup_read_char(__read_only pipe char in_pipe, __global char *dst) 65 ; { 66 ; int gid = get_global_id(0); 67 ; __local reserve_id_t res_id; 68 ; 69 ; res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0)); 70 ; if(is_valid_reserve_id(res_id)) 71 ; { 72 ; read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]); 73 ; work_group_commit_read_pipe(in_pipe, res_id); 74 ; } 75 ; } 76 ; 77 ; __kernel void test_pipe_subgroup_write_uint(__global uint *src, __write_only pipe uint out_pipe) 78 ; { 79 ; int gid = get_global_id(0); 80 ; reserve_id_t res_id; 81 ; 82 ; res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size()); 83 ; if(is_valid_reserve_id(res_id)) 84 ; { 85 ; write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]); 86 ; sub_group_commit_write_pipe(out_pipe, res_id); 87 ; } 88 ; } 89 ; 90 ; __kernel void test_pipe_subgroup_read_uint(__read_only pipe uint in_pipe, __global uint *dst) 91 ; { 92 ; int gid = get_global_id(0); 93 ; reserve_id_t res_id; 94 ; 95 ; res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size()); 96 ; if(is_valid_reserve_id(res_id)) 97 ; { 98 ; read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]); 99 ; sub_group_commit_read_pipe(in_pipe, res_id); 100 ; } 101 ; } 102 103 ; RUN: llvm-as %s -o %t.bc 104 ; RUN: llvm-spirv %t.bc -spirv-text -o %t.spt 105 ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV 106 ; RUN: llvm-spirv %t.bc -o %t.spv 107 ; RUN: llvm-spirv -r %t.spv -o %t.bc 108 ; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM 109 110 ; ModuleID = 'pipe_builtins.cl' 111 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" 112 target triple = "spir64-unknonw-unknown" 113 114 %opencl.reserve_id_t = type opaque 115 %opencl.pipe_t = type opaque 116 117 @test_pipe_workgroup_write_char.res_id = internal unnamed_addr addrspace(3) global %opencl.reserve_id_t* undef, align 8 118 @test_pipe_workgroup_read_char.res_id = internal unnamed_addr addrspace(3) global %opencl.reserve_id_t* undef, align 8 119 120 ; Function Attrs: nounwind 121 define spir_kernel void @test_pipe_convenience_write_uint(i32 addrspace(1)* %src, %opencl.pipe_t addrspace(1)* %out_pipe) #0 { 122 ; CHECK-LLVM-LABEL: @test_pipe_convenience_write_uint 123 ; CHECK-SPIRV-LABEL: 5 Function 124 ; CHECK-SPIRV-NEXT: FunctionParameter 125 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 126 entry: 127 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 128 %sext = shl i64 %call, 32 129 %idxprom = ashr exact i64 %sext, 32 130 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %src, i64 %idxprom 131 %0 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 132 %1 = addrspacecast i8 addrspace(1)* %0 to i8 addrspace(4)* 133 ; CHECK-LLVM: call{{.*}}@_Z10write_pipePU3AS18ocl_pipePU3AS4vjj 134 ; CHECK-SPIRV: WritePipe {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 135 %2 = tail call i32 @_Z10write_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)* %out_pipe, i8 addrspace(4)* %1, i32 4, i32 4) #2 136 ret void 137 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 138 } 139 140 declare spir_func i64 @_Z13get_global_idj(i32) #1 141 142 declare i32 @_Z10write_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)*, i8 addrspace(4)*, i32, i32) 143 144 ; Function Attrs: nounwind 145 define spir_kernel void @test_pipe_convenience_read_uint(%opencl.pipe_t addrspace(1)* %in_pipe, i32 addrspace(1)* %dst) #0 { 146 ; CHECK-LLVM-LABEL: @test_pipe_convenience_read_uint 147 ; CHECK-SPIRV-LABEL: 5 Function 148 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 149 entry: 150 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 151 %sext = shl i64 %call, 32 152 %idxprom = ashr exact i64 %sext, 32 153 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %dst, i64 %idxprom 154 %0 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 155 %1 = addrspacecast i8 addrspace(1)* %0 to i8 addrspace(4)* 156 ; CHECK-LLVM: call{{.*}}@_Z9read_pipePU3AS18ocl_pipePU3AS4vjj 157 ; CHECK-SPIRV: ReadPipe {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 158 %2 = tail call i32 @_Z9read_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)* %in_pipe, i8 addrspace(4)* %1, i32 4, i32 4) #2 159 ret void 160 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 161 } 162 163 declare i32 @_Z9read_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)*, i8 addrspace(4)*, i32, i32) 164 165 ; Function Attrs: nounwind 166 define spir_kernel void @test_pipe_write(i32 addrspace(1)* %src, %opencl.pipe_t addrspace(1)* %out_pipe) #0 { 167 ; CHECK-LLVM-LABEL: @test_pipe_write 168 ; CHECK-SPIRV-LABEL: 5 Function 169 ; CHECK-SPIRV-NEXT: FunctionParameter 170 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 171 entry: 172 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 173 ; CHECK-LLVM: @_Z18reserve_write_pipePU3AS18ocl_pipejjj 174 ; CHECK-SPIRV: ReserveWritePipePackets {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 175 %0 = tail call %opencl.reserve_id_t* @_Z18reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %out_pipe, i32 1, i32 4, i32 4) #2 176 %call1 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %0) #2 177 br i1 %call1, label %if.then, label %if.end 178 179 if.then: ; preds = %entry 180 %sext = shl i64 %call, 32 181 %idxprom = ashr exact i64 %sext, 32 182 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %src, i64 %idxprom 183 %1 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 184 %2 = addrspacecast i8 addrspace(1)* %1 to i8 addrspace(4)* 185 ; CHECK-LLVM: call{{.*}}@_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj 186 ; CHECK-SPIRV: ReservedWritePipe {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 187 %3 = tail call i32 @_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %0, i32 0, i8 addrspace(4)* %2, i32 4, i32 4) #2 188 ; CHECK-LLVM: call{{.*}}@_Z17commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj 189 ; CHECK-SPIRV: CommitWritePipe [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} 190 tail call void @_Z17commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %0, i32 4, i32 4) #2 191 br label %if.end 192 193 if.end: ; preds = %if.then, %entry 194 ret void 195 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 196 } 197 198 declare %opencl.reserve_id_t* @_Z18reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 199 200 declare spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t*) #1 201 202 declare i32 @_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i8 addrspace(4)*, i32, i32) 203 204 declare void @_Z17commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 205 206 ; Function Attrs: nounwind 207 define spir_kernel void @test_pipe_query_functions(%opencl.pipe_t addrspace(1)* %out_pipe, i32 addrspace(1)* nocapture %num_packets, i32 addrspace(1)* nocapture %max_packets) #0 { 208 ; CHECK-LLVM-LABEL: @test_pipe_query_functions 209 ; CHECK-SPIRV-LABEL: 5 Function 210 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 211 entry: 212 ; CHECK-LLVM: call{{.*}}@_Z20get_pipe_max_packetsPU3AS18ocl_pipejj 213 ; CHECK-SPIRV: GetMaxPipePackets {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} 214 %0 = tail call i32 @_Z20get_pipe_max_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)* %out_pipe, i32 4, i32 4) #2 215 store i32 %0, i32 addrspace(1)* %max_packets, align 4, !tbaa !35 216 ; CHECK-LLVM: call{{.*}}@_Z20get_pipe_num_packetsPU3AS18ocl_pipejj 217 ; CHECK-SPIRV: GetNumPipePackets {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} 218 %1 = tail call i32 @_Z20get_pipe_num_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)* %out_pipe, i32 4, i32 4) #2 219 store i32 %1, i32 addrspace(1)* %num_packets, align 4, !tbaa !35 220 ret void 221 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 222 } 223 224 declare i32 @_Z20get_pipe_max_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)*, i32, i32) 225 226 declare i32 @_Z20get_pipe_num_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)*, i32, i32) 227 228 ; Function Attrs: nounwind 229 define spir_kernel void @test_pipe_read(%opencl.pipe_t addrspace(1)* %in_pipe, i32 addrspace(1)* %dst) #0 { 230 ; CHECK-LLVM-LABEL: @test_pipe_read 231 ; CHECK-SPIRV-LABEL: 5 Function 232 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 233 entry: 234 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 235 ; CHECK-LLVM: call{{.*}}@_Z17reserve_read_pipePU3AS18ocl_pipejjj 236 ; CHECK-SPIRV: ReserveReadPipePackets {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 237 %0 = tail call %opencl.reserve_id_t* @_Z17reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %in_pipe, i32 1, i32 4, i32 4) #2 238 %call1 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %0) #2 239 br i1 %call1, label %if.then, label %if.end 240 241 if.then: ; preds = %entry 242 %sext = shl i64 %call, 32 243 %idxprom = ashr exact i64 %sext, 32 244 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %dst, i64 %idxprom 245 %1 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 246 %2 = addrspacecast i8 addrspace(1)* %1 to i8 addrspace(4)* 247 ; CHECK-LLVM: call{{.*}}@_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj 248 ; CHECK-SPIRV: ReservedReadPipe {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 249 %3 = tail call i32 @_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %0, i32 0, i8 addrspace(4)* %2, i32 4, i32 4) #2 250 ; CHECK-LLVM: call{{.*}}@_Z16commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj 251 ; CHECK-SPIRV: CommitReadPipe [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 252 tail call void @_Z16commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %0, i32 4, i32 4) #2 253 br label %if.end 254 255 if.end: ; preds = %if.then, %entry 256 ret void 257 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 258 } 259 260 declare %opencl.reserve_id_t* @_Z17reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 261 262 declare i32 @_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i8 addrspace(4)*, i32, i32) 263 264 declare void @_Z16commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 265 266 ; Function Attrs: nounwind 267 define spir_kernel void @test_pipe_workgroup_write_char(i8 addrspace(1)* %src, %opencl.pipe_t addrspace(1)* %out_pipe) #0 { 268 ; CHECK-LLVM-LABEL: @test_pipe_workgroup_write_char 269 ; CHECK-SPIRV-LABEL: 5 Function 270 ; CHECK-SPIRV-NEXT: FunctionParameter 271 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 272 entry: 273 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 274 %call1 = tail call spir_func i64 @_Z14get_local_sizej(i32 0) #2 275 %0 = trunc i64 %call1 to i32 276 ; CHECK-LLVM: call{{.*}}@_Z29work_group_reserve_write_pipePU3AS18ocl_pipejjj 277 ; CHECK-SPIRV: GroupReserveWritePipePackets {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 278 %1 = tail call %opencl.reserve_id_t* @_Z29work_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %out_pipe, i32 %0, i32 1, i32 1) #2 279 store %opencl.reserve_id_t* %1, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_write_char.res_id, align 8, !tbaa !39 280 %call2 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %1) #2 281 br i1 %call2, label %if.then, label %if.end 282 283 if.then: ; preds = %entry 284 %2 = load %opencl.reserve_id_t*, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_write_char.res_id, align 8, !tbaa !39 285 %call3 = tail call spir_func i64 @_Z12get_local_idj(i32 0) #2 286 %sext = shl i64 %call, 32 287 %idxprom = ashr exact i64 %sext, 32 288 %arrayidx = getelementptr inbounds i8, i8 addrspace(1)* %src, i64 %idxprom 289 %3 = addrspacecast i8 addrspace(1)* %arrayidx to i8 addrspace(4)* 290 %4 = trunc i64 %call3 to i32 291 %5 = tail call i32 @_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %2, i32 %4, i8 addrspace(4)* %3, i32 1, i32 1) #2 292 %6 = load %opencl.reserve_id_t*, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_write_char.res_id, align 8, !tbaa !39 293 ; CHECK-LLVM: call{{.*}}@_Z28work_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj 294 ; CHECK-SPIRV: GroupCommitWritePipe {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 295 tail call void @_Z28work_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %6, i32 1, i32 1) #2 296 br label %if.end 297 298 if.end: ; preds = %if.then, %entry 299 ret void 300 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 301 } 302 303 declare spir_func i64 @_Z14get_local_sizej(i32) #1 304 305 declare %opencl.reserve_id_t* @_Z29work_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 306 307 declare spir_func i64 @_Z12get_local_idj(i32) #1 308 309 declare void @_Z28work_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 310 311 ; Function Attrs: nounwind 312 define spir_kernel void @test_pipe_workgroup_read_char(%opencl.pipe_t addrspace(1)* %in_pipe, i8 addrspace(1)* %dst) #0 { 313 ; CHECK-LLVM-LABEL: @test_pipe_workgroup_read_char 314 ; CHECK-SPIRV-LABEL: 5 Function 315 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 316 entry: 317 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 318 %call1 = tail call spir_func i64 @_Z14get_local_sizej(i32 0) #2 319 %0 = trunc i64 %call1 to i32 320 ; CHECK-LLVM: call{{.*}}@_Z28work_group_reserve_read_pipePU3AS18ocl_pipejjj 321 ; CHECK-SPIRV: GroupReserveReadPipePackets {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 322 %1 = tail call %opencl.reserve_id_t* @_Z28work_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %in_pipe, i32 %0, i32 1, i32 1) #2 323 store %opencl.reserve_id_t* %1, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_read_char.res_id, align 8, !tbaa !39 324 %call2 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %1) #2 325 br i1 %call2, label %if.then, label %if.end 326 327 if.then: ; preds = %entry 328 %2 = load %opencl.reserve_id_t*, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_read_char.res_id, align 8, !tbaa !39 329 %call3 = tail call spir_func i64 @_Z12get_local_idj(i32 0) #2 330 %sext = shl i64 %call, 32 331 %idxprom = ashr exact i64 %sext, 32 332 %arrayidx = getelementptr inbounds i8, i8 addrspace(1)* %dst, i64 %idxprom 333 %3 = addrspacecast i8 addrspace(1)* %arrayidx to i8 addrspace(4)* 334 %4 = trunc i64 %call3 to i32 335 %5 = tail call i32 @_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %2, i32 %4, i8 addrspace(4)* %3, i32 1, i32 1) #2 336 %6 = load %opencl.reserve_id_t*, %opencl.reserve_id_t* addrspace(3)* @test_pipe_workgroup_read_char.res_id, align 8, !tbaa !39 337 ; CHECK-LLVM: call{{.*}}@_Z27work_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj 338 ; CHECK-SPIRV: GroupCommitReadPipe {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 339 tail call void @_Z27work_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %6, i32 1, i32 1) #2 340 br label %if.end 341 342 if.end: ; preds = %if.then, %entry 343 ret void 344 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 345 } 346 347 declare %opencl.reserve_id_t* @_Z28work_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 348 349 declare void @_Z27work_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 350 351 ; Function Attrs: nounwind 352 define spir_kernel void @test_pipe_subgroup_write_uint(i32 addrspace(1)* %src, %opencl.pipe_t addrspace(1)* %out_pipe) #0 { 353 ; CHECK-LLVM-LABEL: @test_pipe_subgroup_write_uint 354 ; CHECK-SPIRV-LABEL: 5 Function 355 ; CHECK-SPIRV-NEXT: FunctionParameter 356 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 357 entry: 358 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 359 %call1 = tail call spir_func i32 @_Z18get_sub_group_sizev() #2 360 ; CHECK-LLVM: call{{.*}}@_Z28sub_group_reserve_write_pipePU3AS18ocl_pipejjj 361 ; CHECK-SPIRV: GroupReserveWritePipePackets {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 362 %0 = tail call %opencl.reserve_id_t* @_Z28sub_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %out_pipe, i32 %call1, i32 4, i32 4) #2 363 %call2 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %0) #2 364 br i1 %call2, label %if.then, label %if.end 365 366 if.then: ; preds = %entry 367 %call3 = tail call spir_func i32 @_Z22get_sub_group_local_idv() #2 368 %sext = shl i64 %call, 32 369 %idxprom = ashr exact i64 %sext, 32 370 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %src, i64 %idxprom 371 %1 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 372 %2 = addrspacecast i8 addrspace(1)* %1 to i8 addrspace(4)* 373 %3 = tail call i32 @_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %0, i32 %call3, i8 addrspace(4)* %2, i32 4, i32 4) #2 374 ; CHECK-LLVM: call{{.*}}@_Z27sub_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj 375 ; CHECK-SPIRV: GroupCommitWritePipe {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 376 tail call void @_Z27sub_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %out_pipe, %opencl.reserve_id_t* %0, i32 4, i32 4) #2 377 br label %if.end 378 379 if.end: ; preds = %if.then, %entry 380 ret void 381 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 382 } 383 384 declare spir_func i32 @_Z18get_sub_group_sizev() #1 385 386 declare %opencl.reserve_id_t* @_Z28sub_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 387 388 declare spir_func i32 @_Z22get_sub_group_local_idv() #1 389 390 declare void @_Z27sub_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 391 392 393 394 ; Function Attrs: nounwind 395 define spir_kernel void @test_pipe_subgroup_read_uint(%opencl.pipe_t addrspace(1)* %in_pipe, i32 addrspace(1)* %dst) #0 { 396 ; CHECK-LLVM-LABEL: @test_pipe_subgroup_read_uint 397 ; CHECK-SPIRV-LABEL: 5 Function 398 ; CHECK-SPIRV-NEXT: FunctionParameter {{[0-9]+}} [[PipeArgID:[0-9]+]] 399 entry: 400 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 401 %call1 = tail call spir_func i32 @_Z18get_sub_group_sizev() #2 402 ; CHECK-LLVM: call{{.*}}@_Z27sub_group_reserve_read_pipePU3AS18ocl_pipejjj 403 ; CHECK-SPIRV: GroupReserveReadPipePackets {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 404 %0 = tail call %opencl.reserve_id_t* @_Z27sub_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)* %in_pipe, i32 %call1, i32 4, i32 4) #2 405 %call2 = tail call spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t* %0) #2 406 br i1 %call2, label %if.then, label %if.end 407 408 if.then: ; preds = %entry 409 %call3 = tail call spir_func i32 @_Z22get_sub_group_local_idv() #2 410 %sext = shl i64 %call, 32 411 %idxprom = ashr exact i64 %sext, 32 412 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %dst, i64 %idxprom 413 %1 = bitcast i32 addrspace(1)* %arrayidx to i8 addrspace(1)* 414 %2 = addrspacecast i8 addrspace(1)* %1 to i8 addrspace(4)* 415 %3 = tail call i32 @_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %0, i32 %call3, i8 addrspace(4)* %2, i32 4, i32 4) #2 416 ; CHECK-LLVM: call{{.*}}@_Z26sub_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj 417 ; CHECK-SPIRV: GroupCommitReadPipe {{[0-9]+}} [[PipeArgID]] {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 418 tail call void @_Z26sub_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)* %in_pipe, %opencl.reserve_id_t* %0, i32 4, i32 4) #2 419 br label %if.end 420 421 if.end: ; preds = %if.then, %entry 422 ret void 423 ; CHECK-SPIRV-LABEL: 1 FunctionEnd 424 } 425 426 declare %opencl.reserve_id_t* @_Z27sub_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 427 428 declare void @_Z26sub_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 429 430 attributes #0 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } 431 attributes #1 = { "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } 432 attributes #2 = { nounwind } 433 434 !opencl.kernels = !{!0, !6, !11, !14, !20, !23, !26, !29, !30} 435 !opencl.enable.FP_CONTRACT = !{} 436 !opencl.spir.version = !{!31} 437 !opencl.ocl.version = !{!32} 438 !opencl.used.extensions = !{!33} 439 !opencl.used.optional.core.features = !{!33} 440 !opencl.compiler.options = !{!33} 441 !llvm.ident = !{!34} 442 443 !0 = !{void (i32 addrspace(1)*, %opencl.pipe_t addrspace(1)*)* @test_pipe_convenience_write_uint, !1, !2, !3, !4, !5} 444 !1 = !{!"kernel_arg_addr_space", i32 1, i32 1} 445 !2 = !{!"kernel_arg_access_qual", !"none", !"write_only"} 446 !3 = !{!"kernel_arg_type", !"uint*", !"uint"} 447 !4 = !{!"kernel_arg_base_type", !"uint*", !"uint"} 448 !5 = !{!"kernel_arg_type_qual", !"", !"pipe"} 449 !6 = !{void (%opencl.pipe_t addrspace(1)*, i32 addrspace(1)*)* @test_pipe_convenience_read_uint, !1, !7, !8, !9, !10} 450 !7 = !{!"kernel_arg_access_qual", !"read_only", !"none"} 451 !8 = !{!"kernel_arg_type", !"uint", !"uint*"} 452 !9 = !{!"kernel_arg_base_type", !"uint", !"uint*"} 453 !10 = !{!"kernel_arg_type_qual", !"pipe", !""} 454 !11 = !{void (i32 addrspace(1)*, %opencl.pipe_t addrspace(1)*)* @test_pipe_write, !1, !2, !12, !13, !5} 455 !12 = !{!"kernel_arg_type", !"int*", !"int"} 456 !13 = !{!"kernel_arg_base_type", !"int*", !"int"} 457 !14 = !{void (%opencl.pipe_t addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*)* @test_pipe_query_functions, !15, !16, !17, !18, !19} 458 !15 = !{!"kernel_arg_addr_space", i32 1, i32 1, i32 1} 459 !16 = !{!"kernel_arg_access_qual", !"write_only", !"none", !"none"} 460 !17 = !{!"kernel_arg_type", !"int", !"int*", !"int*"} 461 !18 = !{!"kernel_arg_base_type", !"int", !"int*", !"int*"} 462 !19 = !{!"kernel_arg_type_qual", !"pipe", !"", !""} 463 !20 = !{void (%opencl.pipe_t addrspace(1)*, i32 addrspace(1)*)* @test_pipe_read, !1, !7, !21, !22, !10} 464 !21 = !{!"kernel_arg_type", !"int", !"int*"} 465 !22 = !{!"kernel_arg_base_type", !"int", !"int*"} 466 !23 = !{void (i8 addrspace(1)*, %opencl.pipe_t addrspace(1)*)* @test_pipe_workgroup_write_char, !1, !2, !24, !25, !5} 467 !24 = !{!"kernel_arg_type", !"char*", !"char"} 468 !25 = !{!"kernel_arg_base_type", !"char*", !"char"} 469 !26 = !{void (%opencl.pipe_t addrspace(1)*, i8 addrspace(1)*)* @test_pipe_workgroup_read_char, !1, !7, !27, !28, !10} 470 !27 = !{!"kernel_arg_type", !"char", !"char*"} 471 !28 = !{!"kernel_arg_base_type", !"char", !"char*"} 472 !29 = !{void (i32 addrspace(1)*, %opencl.pipe_t addrspace(1)*)* @test_pipe_subgroup_write_uint, !1, !2, !3, !4, !5} 473 !30 = !{void (%opencl.pipe_t addrspace(1)*, i32 addrspace(1)*)* @test_pipe_subgroup_read_uint, !1, !7, !8, !9, !10} 474 !31 = !{i32 1, i32 2} 475 !32 = !{i32 2, i32 0} 476 !33 = !{} 477 !34 = !{!"clang version 3.6.1"} 478 !35 = !{!36, !36, i64 0} 479 !36 = !{!"int", !37, i64 0} 480 !37 = !{!"omnipotent char", !38, i64 0} 481 !38 = !{!"Simple C/C++ TBAA"} 482 !39 = !{!40, !40, i64 0} 483 !40 = !{!"reserve_id_t", !37, i64 0} 484
