1;; __kernel void test_32(__global int* res) 2;; { 3;; int tid = get_global_id(0); 4;; 5;; switch(tid) 6;; { 7;; case 0: 8;; res[tid] = 1; 9;; break; 10;; case 1: 11;; res[tid] = 2; 12;; break; 13;; } 14;; } 15;; bash$ clang -cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -include opencl.h -emit-llvm OpSwitch.cl -o test_32.ll 16 17; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV 18 19; CHECK-SPIRV: OpSwitch %[[#]] %[[#]] 0 %[[#]] 1 %[[#]] 20 21define spir_kernel void @test_32(i32 addrspace(1)* %res) { 22entry: 23 %res.addr = alloca i32 addrspace(1)*, align 8 24 %tid = alloca i32, align 4 25 store i32 addrspace(1)* %res, i32 addrspace(1)** %res.addr, align 8 26 %call = call spir_func i64 @_Z13get_global_idj(i32 0) 27 %conv = trunc i64 %call to i32 28 store i32 %conv, i32* %tid, align 4 29 %0 = load i32, i32* %tid, align 4 30 switch i32 %0, label %sw.epilog [ 31 i32 0, label %sw.bb 32 i32 1, label %sw.bb1 33 ] 34 35sw.bb: ; preds = %entry 36 %1 = load i32, i32* %tid, align 4 37 %idxprom = sext i32 %1 to i64 38 %2 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8 39 %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %idxprom 40 store i32 1, i32 addrspace(1)* %arrayidx, align 4 41 br label %sw.epilog 42 43sw.bb1: ; preds = %entry 44 %3 = load i32, i32* %tid, align 4 45 %idxprom2 = sext i32 %3 to i64 46 %4 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8 47 %arrayidx3 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %idxprom2 48 store i32 2, i32 addrspace(1)* %arrayidx3, align 4 49 br label %sw.epilog 50 51sw.epilog: ; preds = %entry, %sw.bb1, %sw.bb 52 ret void 53} 54 55declare spir_func i64 @_Z13get_global_idj(i32) 56