| // 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) nocapture noundef readonly align 8 [[IN:%.*]], ptr addrspace(1) nocapture noundef writeonly align 8 [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] { |
| // CHECK-NEXT: entry: |
| // CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[IN]], i32 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) [[ARRAYIDX1]], align 8, !tbaa [[TBAA7:![0-9]+]] |
| // CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(1) [[OUT]], align 8, !tbaa [[TBAA7]] |
| // 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: [[META3]] = !{i32 1, i32 1} |
| // CHECK: [[META4]] = !{!"none", !"none"} |
| // CHECK: [[META5]] = !{!"long*", !"long*"} |
| // CHECK: [[META6]] = !{!"", !""} |
| // CHECK: [[TBAA7]] = !{[[META8:![0-9]+]], [[META8]], i64 0} |
| // CHECK: [[META8]] = !{!"long", [[META9:![0-9]+]], i64 0} |
| // CHECK: [[META9]] = !{!"omnipotent char", [[META10:![0-9]+]], i64 0} |
| // CHECK: [[META10]] = !{!"Simple C++ TBAA"} |
| //. |