1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4 2//RUN: %clang_cc1 %s -triple spir -emit-llvm -O1 -o - | FileCheck %s 3 4// CHECK-LABEL: define dso_local spir_kernel void @test( 5// CHECK-SAME: ptr addrspace(1) noundef readonly align 8 captures(none) [[IN:%.*]], ptr addrspace(1) noundef writeonly align 8 captures(none) initializes((0, 8)) [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META4:![0-9]+]] !kernel_arg_access_qual [[META5:![0-9]+]] !kernel_arg_type [[META6:![0-9]+]] !kernel_arg_base_type [[META6]] !kernel_arg_type_qual [[META7:![0-9]+]] { 6// CHECK-NEXT: entry: 7// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[IN]], i32 8 8// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) [[ARRAYIDX1]], align 8, !tbaa [[TBAA8:![0-9]+]] 9// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(1) [[OUT]], align 8, !tbaa [[TBAA8]] 10// CHECK-NEXT: ret void 11// 12__kernel void test(__global long *In, __global long *Out) { 13 long m[4] = { In[0], In[1], 0, 0 }; 14 *Out = m[1]; 15} 16//. 17// CHECK: [[META4]] = !{i32 1, i32 1} 18// CHECK: [[META5]] = !{!"none", !"none"} 19// CHECK: [[META6]] = !{!"long*", !"long*"} 20// CHECK: [[META7]] = !{!"", !""} 21// CHECK: [[TBAA8]] = !{[[META9:![0-9]+]], [[META9]], i64 0} 22// CHECK: [[META9]] = !{!"long", [[META10:![0-9]+]], i64 0} 23// CHECK: [[META10]] = !{!"omnipotent char", [[META11:![0-9]+]], i64 0} 24// CHECK: [[META11]] = !{!"Simple C++ TBAA"} 25//. 26