| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 |
| //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 [[META8:![0-9]+]] !kernel_arg_access_qual [[META9:![0-9]+]] !kernel_arg_type [[META10:![0-9]+]] !kernel_arg_base_type [[META10]] !kernel_arg_type_qual [[META11:![0-9]+]] { |
| // CHECK-NEXT: [[ENTRY:.*:]] |
| // CHECK-NEXT: [[ARRAYIDX1_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[IN]], i32 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(1) [[ARRAYIDX1_I]], align 8, !tbaa [[LONG_TBAA12:![0-9]+]] |
| // CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(1) [[OUT]], align 8, !tbaa [[LONG_TBAA12]] |
| // 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: [[META6:![0-9]+]] = !{!"omnipotent char", [[META7:![0-9]+]], i64 0} |
| // CHECK: [[META7]] = !{!"Simple C++ TBAA"} |
| // CHECK: [[META8]] = !{i32 1, i32 1} |
| // CHECK: [[META9]] = !{!"none", !"none"} |
| // CHECK: [[META10]] = !{!"long*", !"long*"} |
| // CHECK: [[META11]] = !{!"", !""} |
| // CHECK: [[LONG_TBAA12]] = !{[[META13:![0-9]+]], [[META13]], i64 0} |
| // CHECK: [[META13]] = !{!"long", [[META6]], i64 0} |
| //. |
