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

   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
  21 define spir_kernel void @test_32(i32 addrspace(1)* %res) {
  22 entry:
  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
  35 sw.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
  43 sw.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
  51 sw.epilog:                                        ; preds = %entry, %sw.bb1, %sw.bb
  52   ret void
  53 }
  54
  55 declare spir_func i64 @_Z13get_global_idj(i32)