Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 8b83ab3

Browse files
fveselovfveselov
authored
[SYCL][ESIMD] noinline functions tests (#143)
1 parent e3c3097 commit 8b83ab3

13 files changed

+477
-0
lines changed

‎SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp

Lines changed: 104 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,104 @@
1+
//===------ noinline_args_size_common.hpp - DPC++ ESIMD on-device test ---===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// The test checks that ESIMD kernels support call of noinline function from
10+
// main function with different total arguments size and retval size. Cases:
11+
// Total arguments size < %arg register size (32 GRFs)
12+
// Total arguments size == %arg register size
13+
// Total arguments size > %arg register size (i.e. stack mem is required)
14+
// Return value size < %retval register size (12 GRFs)
15+
// Return value size == %retval register size
16+
// Return value size > %retval register size
17+
18+
#include "esimd_test_utils.hpp"
19+
20+
#include <CL/sycl.hpp>
21+
#include <CL/sycl/INTEL/esimd.hpp>
22+
#include <iostream>
23+
24+
static_assert(SIZE >= VL, "Size must greater than or equal to VL");
25+
static_assert(SIZE % VL == 0, "Size must be multiple of VL");
26+
constexpr unsigned ROWS = SIZE / VL;
27+
28+
using namespace cl::sycl;
29+
30+
class KernelID;
31+
32+
template <typename TA, typename TB, typename TC>
33+
ESIMD_NOINLINE TC add(TA A, TB B) {
34+
return (TC)A + (TC)B;
35+
}
36+
37+
int main(void) {
38+
queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
39+
40+
auto dev = q.get_device();
41+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
42+
auto ctx = q.get_context();
43+
44+
a_data_t *A = static_cast<a_data_t *>(
45+
sycl::malloc_shared(SIZE * sizeof(a_data_t), dev, ctx));
46+
for (int i = 0; i < SIZE; i++)
47+
A[i] = (a_data_t)1;
48+
49+
b_data_t *B = static_cast<b_data_t *>(
50+
sycl::malloc_shared(SIZE * sizeof(b_data_t), dev, ctx));
51+
for (int i = 0; i < SIZE; i++)
52+
B[i] = (b_data_t)i;
53+
54+
c_data_t *C = static_cast<c_data_t *>(
55+
sycl::malloc_shared(SIZE * sizeof(c_data_t), dev, ctx));
56+
memset(C, 0, SIZE * sizeof(c_data_t));
57+
58+
try {
59+
auto qq = q.submit([&](handler &cgh) {
60+
cgh.parallel_for<KernelID>(
61+
sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
62+
using namespace sycl::INTEL::gpu;
63+
64+
simd<a_data_t, SIZE> va(0);
65+
simd<b_data_t, SIZE> vb(0);
66+
for (int j = 0; j < ROWS; j++) {
67+
va.select<VL, 1>(j * VL) = block_load<a_data_t, VL>(A + j * VL);
68+
vb.select<VL, 1>(j * VL) = block_load<b_data_t, VL>(B + j * VL);
69+
}
70+
71+
auto vc = add<simd<a_data_t, SIZE>, simd<b_data_t, SIZE>,
72+
simd<c_data_t, SIZE>>(va, vb);
73+
74+
for (int j = 0; j < ROWS; j++)
75+
block_store<c_data_t, VL>(C + j * VL, vc.select<VL, 1>(j * VL));
76+
});
77+
});
78+
79+
qq.wait();
80+
} catch (cl::sycl::exception const &e) {
81+
std::cout << "SYCL exception caught: " << e.what() << std::endl;
82+
sycl::free(A, ctx);
83+
sycl::free(B, ctx);
84+
sycl::free(C, ctx);
85+
return e.get_cl_code();
86+
}
87+
88+
unsigned err_cnt = 0;
89+
for (int i = 0; i < SIZE; i++)
90+
if (C[i] != A[i] + B[i])
91+
err_cnt++;
92+
93+
sycl::free(A, ctx);
94+
sycl::free(B, ctx);
95+
sycl::free(C, ctx);
96+
97+
if (err_cnt > 0) {
98+
std::cout << "FAILED" << std::endl;
99+
return 1;
100+
}
101+
102+
std::cout << "passed" << std::endl;
103+
return 0;
104+
}

‎SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_192.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 192;
17+
18+
using a_data_t = int8_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_256.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 256;
17+
18+
using a_data_t = int8_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_512.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 512;
17+
18+
using a_data_t = int8_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===--- noinline_args_char_int_size_64.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 64;
17+
18+
using a_data_t = int8_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===--- noinline_args_char_int_size_96.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 96;
17+
18+
using a_data_t = int8_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_192.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 192;
17+
18+
using a_data_t = int32_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_256.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 256;
17+
18+
using a_data_t = int32_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===-- noinline_args_char_int_size_512.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 512;
17+
18+
using a_data_t = int32_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===--- noinline_args_char_int_size_64.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 64;
17+
18+
using a_data_t = int32_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

‎SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===--- noinline_args_char_int_size_96.cpp - DPC++ ESIMD on-device test -===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda
10+
// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
11+
// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
12+
13+
#include <cstdint>
14+
15+
constexpr unsigned VL = 16;
16+
constexpr unsigned SIZE = 96;
17+
18+
using a_data_t = int32_t;
19+
using b_data_t = int32_t;
20+
using c_data_t = int32_t;
21+
22+
#include "Inputs/noinline_args_size_common.hpp"

0 commit comments

Comments
 (0)