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