| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4 |
| //RUN: %clang_cc1 %s -triple spir -emit-llvm -O1 -o - | FileCheck %s |
| |
| // CHECK-LABEL: define dso_local spir_kernel void @test( |
| // 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]+]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[IN]], i32 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) [[ARRAYIDX1]], align 8, !tbaa [[TBAA8:![0-9]+]] |
| // CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(1) [[OUT]], align 8, !tbaa [[TBAA8]] |
| // CHECK-NEXT: ret void |
| // |
| __kernel void test(__global long *In, __global long *Out) { |
| long m[4] = { In[0], In[1], 0, 0 }; |
| *Out = m[1]; |
| } |
| //. |
| // CHECK: [[META4]] = !{i32 1, i32 1} |
| // CHECK: [[META5]] = !{!"none", !"none"} |
| // CHECK: [[META6]] = !{!"long*", !"long*"} |
| // CHECK: [[META7]] = !{!"", !""} |
| // CHECK: [[TBAA8]] = !{[[META9:![0-9]+]], [[META9]], i64 0} |
| // CHECK: [[META9]] = !{!"long", [[META10:![0-9]+]], i64 0} |
| // CHECK: [[META10]] = !{!"omnipotent char", [[META11:![0-9]+]], i64 0} |
| // CHECK: [[META11]] = !{!"Simple C++ TBAA"} |
| //. |
