Skip to content

Commit 219439c

Browse files
KornevNikitasarnex
KornevNikita
and
sarnex
authored
[CI][sycl-rel] Upd workflow (#19293)
Aligned with regular nightly. Removed aws-cuda job, I believe it's not required here. Also updated tests failed on DG2 due to old driver. --------- Co-authored-by: Nick Sarnie <[email protected]>
1 parent 3531dcf commit 219439c

File tree

7 files changed

+50
-64
lines changed

7 files changed

+50
-64
lines changed

‎.github/workflows/sycl-rel-nightly.yml

Lines changed: 10 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,11 @@ jobs:
3232
image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
3333
target_devices: hip:gpu
3434

35+
- name: NVIDIA/CUDA
36+
runner: '["Linux", "cuda"]'
37+
image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
38+
target_devices: cuda:gpu
39+
3540
- name: Intel L0 Gen12 GPU
3641
runner: '["Linux", "gen12"]'
3742
image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
@@ -43,6 +48,11 @@ jobs:
4348
image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
4449
target_devices: level_zero:gpu
4550

51+
- name: Intel L0 Arc A-Series GPU
52+
runner: '["Linux", "arc"]'
53+
image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
54+
target_devices: level_zero:gpu
55+
4656
- name: Intel OCL Gen12 GPU
4757
runner: '["Linux", "gen12"]'
4858
image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
@@ -102,40 +112,6 @@ jobs:
102112
extra_lit_opts: ${{ matrix.extra_lit_opts }}
103113
repo_ref: ${{ github.sha }}
104114

105-
cuda-aws-start:
106-
needs: [ubuntu2204_build]
107-
if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}
108-
uses: ./.github/workflows/sycl-aws.yml
109-
secrets: inherit
110-
with:
111-
mode: start
112-
ref: ${{ github.sha }}
113-
114-
cuda-run-tests:
115-
needs: [ubuntu2204_build, cuda-aws-start]
116-
if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}
117-
uses: ./.github/workflows/sycl-linux-run-tests.yml
118-
with:
119-
name: CUDA E2E
120-
runner: '["aws_cuda-${{ github.run_id }}-${{ github.run_attempt }}"]'
121-
image: ghcr.io/intel/llvm/ubuntu2204_build:latest-133fee559371ce0e6ff867e378c21cde2bdf6c90
122-
image_options: -u 1001 --gpus all --cap-add SYS_ADMIN --env NVIDIA_DISABLE_REQUIRE=1
123-
target_devices: cuda:gpu
124-
repo_ref: ${{ github.sha }}
125-
126-
sycl_toolchain_artifact: sycl_linux_default
127-
sycl_toolchain_archive: ${{ needs.ubuntu2204_build.outputs.artifact_archive_name }}
128-
sycl_toolchain_decompress_command: ${{ needs.ubuntu2204_build.outputs.artifact_decompress_command }}
129-
130-
cuda-aws-stop:
131-
needs: [cuda-aws-start, cuda-run-tests]
132-
if: always() && ${{ needs.cuda-aws-start.result != 'skipped' }}
133-
uses: ./.github/workflows/sycl-aws.yml
134-
secrets: inherit
135-
with:
136-
mode: stop
137-
ref: ${{ github.sha }}
138-
139115
build-sycl-cts:
140116
needs: ubuntu2204_build
141117
if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}

‎sycl/test-e2e/ESIMD/group_barrier.cpp

Lines changed: 37 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -14,40 +14,54 @@
1414
#include "esimd_test_utils.hpp"
1515
#include <sycl/ext/oneapi/experimental/root_group.hpp>
1616
#include <sycl/group_barrier.hpp>
17+
#include <sycl/kernel_bundle.hpp>
1718

18-
static constexpr int WorkGroupSize = 16;
19+
namespace syclex = sycl::ext::oneapi::experimental;
20+
21+
static constexpr int WorkGroupSize = 32;
1922

2023
static constexpr int VL = 16;
24+
25+
template <int Val> class MyKernel;
26+
2127
template <bool UseThisWorkItemAPI> bool test(sycl::queue &q) {
2228
bool Pass = true;
23-
const auto MaxWGs = 8;
24-
size_t WorkItemCount = MaxWGs * WorkGroupSize * VL;
2529
std::cout << "Test case UseThisWorkItemAPI="
2630
<< std::to_string(UseThisWorkItemAPI) << std::endl;
2731
const auto Props = sycl::ext::oneapi::experimental::properties{
2832
sycl::ext::oneapi::experimental::use_root_sync};
29-
sycl::buffer<int> DataBuf{sycl::range{WorkItemCount}};
30-
const auto Range = sycl::nd_range<1>{MaxWGs * WorkGroupSize, WorkGroupSize};
33+
auto Bundle =
34+
sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context());
35+
auto Kernel = Bundle.template get_kernel<MyKernel<UseThisWorkItemAPI>>();
36+
sycl::range<3> LocalRange{WorkGroupSize, 1, 1};
37+
auto MaxWGs = Kernel.template ext_oneapi_get_info<
38+
syclex::info::kernel_queue_specific::max_num_work_groups>(q, LocalRange,
39+
0);
40+
auto GlobalRange = LocalRange;
41+
size_t WorkItemCount = GlobalRange.size() * VL;
42+
sycl::buffer<int> DataBuf{WorkItemCount};
43+
const auto Range = sycl::nd_range<3>{GlobalRange, LocalRange};
3144
q.submit([&](sycl::handler &h) {
3245
sycl::accessor Data{DataBuf, h};
33-
h.parallel_for(Range, Props, [=](sycl::nd_item<1> it) SYCL_ESIMD_KERNEL {
34-
int ID = it.get_global_linear_id();
35-
__ESIMD_NS::simd<int, VL> V(ID, 1);
36-
// Write data to another kernel's data to verify the barrier works.
37-
__ESIMD_NS::block_store(
38-
Data, (WorkItemCount * sizeof(int)) - (ID * sizeof(int) * VL), V);
39-
if constexpr (UseThisWorkItemAPI) {
40-
auto Root =
41-
sycl::ext::oneapi::experimental::this_work_item::get_root_group<
42-
1>();
43-
sycl::group_barrier(Root);
44-
} else {
45-
auto Root = it.ext_oneapi_get_root_group();
46-
sycl::group_barrier(Root);
47-
}
48-
__ESIMD_NS::simd<int, VL> VOther(ID * VL, 1);
49-
__ESIMD_NS::block_store(Data, ID * sizeof(int) * VL, VOther);
50-
});
46+
h.parallel_for<MyKernel<UseThisWorkItemAPI>>(
47+
Range, Props, [=](sycl::nd_item<3> it) SYCL_ESIMD_KERNEL {
48+
int ID = it.get_global_linear_id();
49+
__ESIMD_NS::simd<int, VL> V(ID, 1);
50+
// Write data to another kernel's data to verify the barrier works.
51+
__ESIMD_NS::block_store(
52+
Data, (WorkItemCount * sizeof(int)) - (ID * sizeof(int) * VL),
53+
V);
54+
if constexpr (UseThisWorkItemAPI) {
55+
auto Root = sycl::ext::oneapi::experimental::this_work_item::
56+
get_root_group<1>();
57+
sycl::group_barrier(Root);
58+
} else {
59+
auto Root = it.ext_oneapi_get_root_group();
60+
sycl::group_barrier(Root);
61+
}
62+
__ESIMD_NS::simd<int, VL> VOther(ID * VL, 1);
63+
__ESIMD_NS::block_store(Data, ID * sizeof(int) * VL, VOther);
64+
});
5165
}).wait();
5266
sycl::host_accessor Data{DataBuf};
5367
int ErrCnt = 0;

‎sycl/test-e2e/Graph/Explicit/buffer_ordering.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
// XFAIL: run-mode && gpu-intel-dg2
2-
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18579
31
// RUN: %{build} -o %t.out
42
// RUN: %{run} %t.out
53
// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG

‎sycl/test-e2e/Graph/RecordReplay/buffer_ordering.cpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
// XFAIL: run-mode && gpu-intel-dg2
2-
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18579
31
// RUN: %{build} -o %t.out
42
// RUN: %{run} %t.out
53
// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG

‎sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_SLM.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
// UNSUPPORTED: target-nvidia, target-amd
99
// UNSUPPORTED-INTENDED: aspect-ext_intel_matrix isn't currently supported for
1010
// other triples
11-
// XFAIL: run-mode && gpu-intel-dg2
11+
// XFAIL: run-mode && igc-dev
1212
// XFAIL-TRACKER: CMPLRLLVM-66371
1313

1414
// REQUIRES: aspect-ext_intel_matrix, gpu

‎sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_arg_dim.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717

1818
// Waiting for the commit in IGC to be pulled into the driver to resolve the
1919
// test.
20-
// XFAIL: gpu-intel-dg2 && run-mode
20+
// XFAIL: run-mode && igc-dev
2121
// XFAIL-TRACKER: GSD-10510
2222

2323
#include "common.hpp"

‎sycl/test-e2e/Matrix/joint_matrix_bf16_fill_k_cache_runtime_dim.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818

1919
// Waiting for the commit in IGC to be pulled into the driver to resolve the
2020
// test.
21-
// XFAIL: gpu-intel-dg2 && run-mode
21+
// XFAIL: run-mode && igc-dev
2222
// XFAIL-TRACKER: GSD-10510
2323

2424
#include "common.hpp"

0 commit comments

Comments
 (0)