llvm/test/CodeGen/SPIRV/branching/OpSwitch64.ll

   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
  24 define spir_kernel void @test_64(i32 addrspace(1)* %res) {
  25 entry:
  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
  38 sw.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
  45 sw.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
  52 sw.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
  59 sw.epilog:                                        ; preds = %entry, %sw.bb3, %sw.bb1, %sw.bb
  60   ret void
  61 }
  62
  63 declare spir_func i64 @_Z13get_global_idj(i32)