Skip to content

[SYCL] Add a test for generated device IR for sycl::vec::as() #13750

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
aelovikov-intel merged 2 commits into from
May 10, 2024
Merged

[SYCL] Add a test for generated device IR for sycl::vec::as() #13750

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
2b841c7
[SYCL] Add a test for generated device IR for `sycl::vec::as()`
aelovikov-intel May 10, 2024
a3e39b0
Limit to Linux only
aelovikov-intel May 10, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
54 changes: 54 additions & 0 deletions sycl/test/check_device_code/vector/vector_as.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Copy link
Contributor Author

@aelovikov-intel aelovikov-intel May 10, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-format complains here but that line has to stay as the script needs it like this. Normally it's used under clang/test and long lines are allowed there.

// 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 <sycl/sycl.hpp>

template SYCL_EXTERNAL sycl::vec<int, 4> sycl::vec<float, 4>::as<sycl::vec<int, 4>>() 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