llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll

   1 ; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
   2
   3 ;; The IR was generated from the following source:
   4 ;; #include <CL/sycl.hpp>
   5 ;;
   6 ;; int main() {
   7 ;;   sycl::queue Queue;
   8 ;;   int array[2][3] = {0};
   9 ;;   {
  10 ;;     sycl::range<2> Range(2, 3);
  11 ;;     sycl::buffer<int, 2> buf((int *)array, Range,
  12 ;;                              {cl::sycl::property::buffer::use_host_ptr()});
  13 ;;
  14 ;;     Queue.submit([&](sycl::handler &cgh) {
  15 ;;       auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
  16 ;;       cgh.parallel_for<class dim2_subscr>(Range, [=](sycl::item<2> itemID) {
  17 ;;         acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id();
  18 ;;       });
  19 ;;     });
  20 ;;     Queue.wait();
  21 ;;   }
  22 ;;   return 0;
  23 ;; }
  24 ;; Command line:
  25 ;; clang++ -fsycl -fsycl-device-only emit-llvm tmp.cpp -o tmp.bc
  26 ;; llvm-spirv tmp.bc -spirv-text -o builtin_vars_arithmetics.ll
  27
  28 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId:]] BuiltIn GlobalInvocationId
  29 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize:]] BuiltIn GlobalSize
  30 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset:]] BuiltIn GlobalOffset
  31 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId]] Constant
  32 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize]] Constant
  33 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset]] Constant
  34 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset]] LinkageAttributes "__spirv_BuiltInGlobalOffset" Import
  35 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize]] LinkageAttributes "__spirv_BuiltInGlobalSize" Import
  36 ; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId]] LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
  37
  38 %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" }
  39 %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" = type { [2 x i64] }
  40 %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" }
  41
  42 $"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11dim2_subscr" = comdat any
  43
  44 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
  45 @__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
  46 @__spirv_BuiltInGlobalOffset = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
  47
  48 define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11dim2_subscr"(i32 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr comdat {
  49 entry:
  50   %agg.tmp4.sroa.0.sroa.2.0.agg.tmp4.sroa.0.0..sroa_cast.sroa_idx65 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range"* %_arg_2, i64 0, i32 0, i32 0, i64 1
  51   %agg.tmp4.sroa.0.sroa.2.0.copyload = load i64, i64* %agg.tmp4.sroa.0.sroa.2.0.agg.tmp4.sroa.0.0..sroa_cast.sroa_idx65, align 8
  52   %agg.tmp5.sroa.0.sroa.0.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
  53   %agg.tmp5.sroa.0.sroa.0.0.copyload = load i64, i64* %agg.tmp5.sroa.0.sroa.0.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx, align 8
  54   %agg.tmp5.sroa.0.sroa.2.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx69 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 1
  55   %agg.tmp5.sroa.0.sroa.2.0.copyload = load i64, i64* %agg.tmp5.sroa.0.sroa.2.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx69, align 8
  56   %0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
  57   %1 = extractelement <3 x i64> %0, i64 1
  58   %2 = extractelement <3 x i64> %0, i64 0
  59   %3 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalSize to <3 x i64> addrspace(4)*), align 32
  60   %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalOffset to <3 x i64> addrspace(4)*), align 32
  61   %5 = sub <3 x i64> %0, %4
  62   %6 = sub <3 x i64> %0, %4
  63   %7 = extractelement <3 x i64> %6, i64 0
  64   %8 = extractelement <3 x i64> %5, i32 1
  65   %9 = extractelement <3 x i64> %3, i64 0
  66   %10 = mul i64 %8, %9
  67   %add.i.i.i = add i64 %7, %10
  68   %add6.i.i.i.i = add i64 %1, %agg.tmp5.sroa.0.sroa.0.0.copyload
  69   %mul.1.i.i.i.i = mul i64 %add6.i.i.i.i, %agg.tmp4.sroa.0.sroa.2.0.copyload
  70   %add.1.i.i.i.i = add i64 %2, %agg.tmp5.sroa.0.sroa.2.0.copyload
  71   %add6.1.i.i.i.i = add i64 %add.1.i.i.i.i, %mul.1.i.i.i.i
  72   %ptridx.i.i.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %add6.1.i.i.i.i
  73   %ptridx.ascast.i.i.i = addrspacecast i32 addrspace(1)* %ptridx.i.i.i to i32 addrspace(4)*
  74   %11 = load i32, i32 addrspace(4)* %ptridx.ascast.i.i.i, align 4
  75   %12 = trunc i64 %add.i.i.i to i32
  76   %conv5.i = add i32 %11, %12
  77   store i32 %conv5.i, i32 addrspace(4)* %ptridx.ascast.i.i.i, align 4
  78   ret void
  79 }