81 lines · plain
1; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV2; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}3 4;; The IR was generated from the following source:5;; #include <CL/sycl.hpp>6;;7;; int main() {8;; sycl::queue Queue;9;; int array[2][3] = {0};10;; {11;; sycl::range<2> Range(2, 3);12;; sycl::buffer<int, 2> buf((int *)array, Range,13;; {cl::sycl::property::buffer::use_host_ptr()});14;;15;; Queue.submit([&](sycl::handler &cgh) {16;; auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);17;; cgh.parallel_for<class dim2_subscr>(Range, [=](sycl::item<2> itemID) {18;; acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id();19;; });20;; });21;; Queue.wait();22;; }23;; return 0;24;; }25;; Command line:26;; clang++ -fsycl -fsycl-device-only emit-llvm tmp.cpp -o tmp.bc27;; llvm-spirv tmp.bc -spirv-text -o builtin_vars_arithmetics.ll28 29; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId:]] BuiltIn GlobalInvocationId30; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize:]] BuiltIn GlobalSize31; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset:]] BuiltIn GlobalOffset32; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId]] Constant33; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize]] Constant34; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset]] Constant35; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalOffset]] LinkageAttributes "__spirv_BuiltInGlobalOffset" Import36; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalSize]] LinkageAttributes "__spirv_BuiltInGlobalSize" Import37; CHECK-SPIRV-DAG: OpDecorate %[[#GlobalInvocationId]] LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import38 39%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" }40%"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" = type { [2 x i64] }41%"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" }42 43$"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE11dim2_subscr" = comdat any44 45@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3246@__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3247@__spirv_BuiltInGlobalOffset = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3248 49define 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 {50entry:51 %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 152 %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 853 %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 054 %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 855 %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 156 %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 857 %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 3258 %1 = extractelement <3 x i64> %0, i64 159 %2 = extractelement <3 x i64> %0, i64 060 %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 3261 %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 3262 %5 = sub <3 x i64> %0, %463 %6 = sub <3 x i64> %0, %464 %7 = extractelement <3 x i64> %6, i64 065 %8 = extractelement <3 x i64> %5, i32 166 %9 = extractelement <3 x i64> %3, i64 067 %10 = mul i64 %8, %968 %add.i.i.i = add i64 %7, %1069 %add6.i.i.i.i = add i64 %1, %agg.tmp5.sroa.0.sroa.0.0.copyload70 %mul.1.i.i.i.i = mul i64 %add6.i.i.i.i, %agg.tmp4.sroa.0.sroa.2.0.copyload71 %add.1.i.i.i.i = add i64 %2, %agg.tmp5.sroa.0.sroa.2.0.copyload72 %add6.1.i.i.i.i = add i64 %add.1.i.i.i.i, %mul.1.i.i.i.i73 %ptridx.i.i.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %add6.1.i.i.i.i74 %ptridx.ascast.i.i.i = addrspacecast i32 addrspace(1)* %ptridx.i.i.i to i32 addrspace(4)*75 %11 = load i32, i32 addrspace(4)* %ptridx.ascast.i.i.i, align 476 %12 = trunc i64 %add.i.i.i to i3277 %conv5.i = add i32 %11, %1278 store i32 %conv5.i, i32 addrspace(4)* %ptridx.ascast.i.i.i, align 479 ret void80}81