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' 111target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" 112target 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 121define 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]+]] 126entry: 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 140declare spir_func i64 @_Z13get_global_idj(i32) #1 141 142declare i32 @_Z10write_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)*, i8 addrspace(4)*, i32, i32) 143 144; Function Attrs: nounwind 145define 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]+]] 149entry: 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 163declare i32 @_Z9read_pipePU3AS18ocl_pipePU3AS4vjj(%opencl.pipe_t addrspace(1)*, i8 addrspace(4)*, i32, i32) 164 165; Function Attrs: nounwind 166define 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]+]] 171entry: 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 179if.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 193if.end: ; preds = %if.then, %entry 194 ret void 195; CHECK-SPIRV-LABEL: 1 FunctionEnd 196} 197 198declare %opencl.reserve_id_t* @_Z18reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 199 200declare spir_func zeroext i1 @_Z19is_valid_reserve_id13ocl_reserveid(%opencl.reserve_id_t*) #1 201 202declare i32 @_Z10write_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i8 addrspace(4)*, i32, i32) 203 204declare void @_Z17commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 205 206; Function Attrs: nounwind 207define 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]+]] 211entry: 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 224declare i32 @_Z20get_pipe_max_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)*, i32, i32) 225 226declare i32 @_Z20get_pipe_num_packetsPU3AS18ocl_pipejj(%opencl.pipe_t addrspace(1)*, i32, i32) 227 228; Function Attrs: nounwind 229define 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]+]] 233entry: 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 241if.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 255if.end: ; preds = %if.then, %entry 256 ret void 257; CHECK-SPIRV-LABEL: 1 FunctionEnd 258} 259 260declare %opencl.reserve_id_t* @_Z17reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 261 262declare i32 @_Z9read_pipePU3AS18ocl_pipe13ocl_reserveidjPU3AS4vjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i8 addrspace(4)*, i32, i32) 263 264declare void @_Z16commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 265 266; Function Attrs: nounwind 267define 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]+]] 272entry: 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 283if.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 298if.end: ; preds = %if.then, %entry 299 ret void 300; CHECK-SPIRV-LABEL: 1 FunctionEnd 301} 302 303declare spir_func i64 @_Z14get_local_sizej(i32) #1 304 305declare %opencl.reserve_id_t* @_Z29work_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 306 307declare spir_func i64 @_Z12get_local_idj(i32) #1 308 309declare void @_Z28work_group_commit_write_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 310 311; Function Attrs: nounwind 312define 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]+]] 316entry: 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 327if.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 342if.end: ; preds = %if.then, %entry 343 ret void 344; CHECK-SPIRV-LABEL: 1 FunctionEnd 345} 346 347declare %opencl.reserve_id_t* @_Z28work_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 348 349declare void @_Z27work_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 350 351; Function Attrs: nounwind 352define 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]+]] 357entry: 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 366if.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 379if.end: ; preds = %if.then, %entry 380 ret void 381; CHECK-SPIRV-LABEL: 1 FunctionEnd 382} 383 384declare spir_func i32 @_Z18get_sub_group_sizev() #1 385 386declare %opencl.reserve_id_t* @_Z28sub_group_reserve_write_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 387 388declare spir_func i32 @_Z22get_sub_group_local_idv() #1 389 390declare 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 395define 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]+]] 399entry: 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 408if.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 421if.end: ; preds = %if.then, %entry 422 ret void 423; CHECK-SPIRV-LABEL: 1 FunctionEnd 424} 425 426declare %opencl.reserve_id_t* @_Z27sub_group_reserve_read_pipePU3AS18ocl_pipejjj(%opencl.pipe_t addrspace(1)*, i32, i32, i32) 427 428declare void @_Z26sub_group_commit_read_pipePU3AS18ocl_pipe13ocl_reserveidjj(%opencl.pipe_t addrspace(1)*, %opencl.reserve_id_t*, i32, i32) 429 430attributes #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" } 431attributes #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" } 432attributes #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