diff --git a/sycl/test/check_device_code/vector/vector_as.cpp b/sycl/test/check_device_code/vector/vector_as.cpp new file mode 100644 index 0000000000000..1079d4658f9ab --- /dev/null +++ b/sycl/test/check_device_code/vector/vector_as.cpp @@ -0,0 +1,54 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions "as" --include-generated-funcs --version 4 +// NOTE: "%if preview-breaking-changes-supported" guard has to be temporarily +// NOTE: removed/disabled to re-generate the checks. + +// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s +// RUN: %if preview-breaking-changes-supported %{ \ +// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s -fpreview-breaking-changes | FileCheck %s --check-prefix=CHECK-PREVIEW \ +// RUN: %} + +// Windows/linux have some slight differences in IR generation (function +// arguments passing and long/long long differences/mangling) that could +// complicate test updates while not improving test coverage. Limiting to linux +// should be fine. +// REQUIRES: linux + +#include + +template SYCL_EXTERNAL sycl::vec sycl::vec::as>() const; +// CHECK-LABEL: define weak_odr dso_local spir_func void @_ZNK4sycl3_V13vecIfLi4EE2asINS1_IiLi4EEEEET_v( +// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: br label [[FOR_COND_I:%.*]] +// CHECK: for.cond.i: +// CHECK-NEXT: [[I_0_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I:%.*]], [[FOR_BODY_I:%.*]] ] +// CHECK-NEXT: [[CMP_I:%.*]] = icmp ult i64 [[I_0_I]], 16 +// CHECK-NEXT: br i1 [[CMP_I]], label [[FOR_BODY_I]], label [[_ZN4SYCL3_V16DETAIL6MEMCPYEPVPKVM_EXIT:%.*]] +// CHECK: for.body.i: +// CHECK-NEXT: [[ARRAYIDX_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[THIS]], i64 [[I_0_I]] +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[ARRAYIDX_I]], align 1, !tbaa [[TBAA7:![0-9]+]] +// CHECK-NEXT: [[ARRAYIDX1_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[AGG_RESULT]], i64 [[I_0_I]] +// CHECK-NEXT: store i8 [[TMP0]], ptr addrspace(4) [[ARRAYIDX1_I]], align 1, !tbaa [[TBAA7]] +// CHECK-NEXT: [[INC_I]] = add nuw nsw i64 [[I_0_I]], 1 +// CHECK-NEXT: br label [[FOR_COND_I]], !llvm.loop [[LOOP10:![0-9]+]] +// CHECK: _ZN4sycl3_V16detail6memcpyEPvPKvm.exit: +// CHECK-NEXT: ret void +// +// +// CHECK-PREVIEW-LABEL: define weak_odr dso_local spir_func void @_ZNK4sycl3_V13vecIfLi4EE2asINS1_IiLi4EEEEET_v( +// CHECK-PREVIEW-SAME: ptr addrspace(4) dead_on_unwind noalias writable sret(%"class.sycl::_V1::vec") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef align 16 dereferenceable_or_null(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] comdat align 2 !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] { +// CHECK-PREVIEW-NEXT: entry: +// CHECK-PREVIEW-NEXT: br label [[FOR_COND_I:%.*]] +// CHECK-PREVIEW: for.cond.i: +// CHECK-PREVIEW-NEXT: [[I_0_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ [[INC_I:%.*]], [[FOR_BODY_I:%.*]] ] +// CHECK-PREVIEW-NEXT: [[CMP_I:%.*]] = icmp ult i64 [[I_0_I]], 16 +// CHECK-PREVIEW-NEXT: br i1 [[CMP_I]], label [[FOR_BODY_I]], label [[_ZN4SYCL3_V16DETAIL6MEMCPYEPVPKVM_EXIT:%.*]] +// CHECK-PREVIEW: for.body.i: +// CHECK-PREVIEW-NEXT: [[ARRAYIDX_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[THIS]], i64 [[I_0_I]] +// CHECK-PREVIEW-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[ARRAYIDX_I]], align 1, !tbaa [[TBAA7:![0-9]+]] +// CHECK-PREVIEW-NEXT: [[ARRAYIDX1_I:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[AGG_RESULT]], i64 [[I_0_I]] +// CHECK-PREVIEW-NEXT: store i8 [[TMP0]], ptr addrspace(4) [[ARRAYIDX1_I]], align 1, !tbaa [[TBAA7]] +// CHECK-PREVIEW-NEXT: [[INC_I]] = add nuw nsw i64 [[I_0_I]], 1 +// CHECK-PREVIEW-NEXT: br label [[FOR_COND_I]], !llvm.loop [[LOOP10:![0-9]+]] +// CHECK-PREVIEW: _ZN4sycl3_V16detail6memcpyEPvPKvm.exit: +// CHECK-PREVIEW-NEXT: ret void