Skip to content

Commit 56bbbfa

Browse files
jbrodmanjbrodman
committed
Move spec constants and printf to ONEAPI::experimental
Signed-off-by: James Brodman <[email protected]>
1 parent e0b8033 commit 56bbbfa

File tree

17 files changed

+82
-74
lines changed

17 files changed

+82
-74
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2813,10 +2813,11 @@ bool Util::isSyclHalfType(const QualType &Ty) {
28132813

28142814
bool Util::isSyclSpecConstantType(const QualType &Ty) {
28152815
const StringRef &Name = "spec_constant";
2816-
std::array<DeclContextDesc, 4> Scopes = {
2816+
std::array<DeclContextDesc, 5> Scopes = {
28172817
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
28182818
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
28192819
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"},
2820+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "experimental"},
28202821
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
28212822
return matchQualifiedTypeName(Ty, Scopes);
28222823
}

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -243,6 +243,7 @@ struct get_kernel_name_t<auto_name, Type> {
243243
};
244244

245245
namespace ONEAPI {
246+
namespace experimental {
246247
template <typename T, typename ID = T>
247248
class spec_constant {
248249
public:
@@ -256,6 +257,7 @@ class spec_constant {
256257
return get();
257258
}
258259
};
260+
} // namespace experimental
259261
} // namespace ONEAPI
260262

261263
#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))

‎clang/test/CodeGenSYCL/int_header_spec_const.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -20,18 +20,18 @@ class MyDoubleConst;
2020

2121
int main() {
2222
// Create specialization constants.
23-
cl::sycl::ONEAPI::spec_constant<bool, MyBoolConst> i1(false);
24-
cl::sycl::ONEAPI::spec_constant<char, MyInt8Const> i8(0);
25-
cl::sycl::ONEAPI::spec_constant<unsigned char, MyUInt8Const> ui8(0);
26-
cl::sycl::ONEAPI::spec_constant<short, MyInt16Const> i16(0);
27-
cl::sycl::ONEAPI::spec_constant<unsigned short, MyUInt16Const> ui16(0);
28-
cl::sycl::ONEAPI::spec_constant<int, MyInt32Const> i32(0);
23+
cl::sycl::ONEAPI::experimental::spec_constant<bool, MyBoolConst> i1(false);
24+
cl::sycl::ONEAPI::experimental::spec_constant<char, MyInt8Const> i8(0);
25+
cl::sycl::ONEAPI::experimental::spec_constant<unsigned char, MyUInt8Const> ui8(0);
26+
cl::sycl::ONEAPI::experimental::spec_constant<short, MyInt16Const> i16(0);
27+
cl::sycl::ONEAPI::experimental::spec_constant<unsigned short, MyUInt16Const> ui16(0);
28+
cl::sycl::ONEAPI::experimental::spec_constant<int, MyInt32Const> i32(0);
2929
// Constant used twice, but there must be single entry in the int header,
3030
// otherwise compilation error would be issued.
31-
cl::sycl::ONEAPI::spec_constant<int, MyInt32Const> i32_1(0);
32-
cl::sycl::ONEAPI::spec_constant<unsigned int, MyUInt32Const> ui32(0);
33-
cl::sycl::ONEAPI::spec_constant<float, MyFloatConst> f32(0);
34-
cl::sycl::ONEAPI::spec_constant<double, MyDoubleConst> f64(0);
31+
cl::sycl::ONEAPI::experimental::spec_constant<int, MyInt32Const> i32_1(0);
32+
cl::sycl::ONEAPI::experimental::spec_constant<unsigned int, MyUInt32Const> ui32(0);
33+
cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32(0);
34+
cl::sycl::ONEAPI::experimental::spec_constant<double, MyDoubleConst> f64(0);
3535

3636
double val;
3737
double *ptr = &val; // to avoid "unused" warnings

‎clang/test/SemaSYCL/Inputs/sycl.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,11 +209,11 @@ class handler {
209209
};
210210

211211
namespace ONEAPI {
212-
212+
namespace experimental {
213213
template <typename T, typename ID = T>
214214
class spec_constant {};
215215
} // namespace experimental
216-
216+
} // namespace ONEAPI
217217
} // namespace sycl
218218
} // namespace cl
219219

‎clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,9 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
1010
}
1111

1212
int main() {
13-
cl::sycl::ONEAPI::spec_constant<char, class MyInt32Const> spec_const;
13+
cl::sycl::ONEAPI::experimental::spec_constant<char, class MyInt32Const> spec_const;
1414
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessor;
15-
// CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::ONEAPI::spec_constant<char, class MyInt32Const>'
15+
// CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::ONEAPI::experimental::spec_constant<char, class MyInt32Const>'
1616
// CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>'
1717
kernel<class MyKernel>([spec_const, accessor]() {});
1818
return 0;

‎llvm/test/tools/sycl-post-link/sc_sym_two_refs.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
88
target triple = "spir64-unknown-unknown-sycldevice"
99

10-
%"sycl::experimental::spec_constant" = type { i8 }
10+
%"sycl::ONEAPI::experimental::spec_constant" = type { i8 }
1111

1212
@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
1313
; CHECK-NOT: @SCSymID
@@ -21,7 +21,7 @@ define weak_odr dso_local spir_kernel void @Kernel() {
2121
}
2222

2323
; Function Attrs: norecurse
24-
define dso_local spir_func float @foo_float(%"sycl::experimental::spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
24+
define dso_local spir_func float @foo_float(%"sycl::ONEAPI::experimental::spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
2525
%2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
2626
ret float %2
2727
}

‎sycl/include/CL/sycl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@
2828
#include <CL/sycl/nd_item.hpp>
2929
#include <CL/sycl/nd_range.hpp>
3030
#include <CL/sycl/ONEAPI/atomic.hpp>
31-
#include <CL/sycl/ONEAPI/builtins.hpp>
31+
#include <CL/sycl/ONEAPI/experimental/builtins.hpp>
3232
#include <CL/sycl/ONEAPI/function_pointer.hpp>
3333
#include <CL/sycl/ONEAPI/group_algorithm.hpp>
3434
#include <CL/sycl/ONEAPI/reduction.hpp>

‎sycl/include/CL/sycl/ONEAPI/builtins.hpp renamed to ‎sycl/include/CL/sycl/ONEAPI/experimental/builtins.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
__SYCL_INLINE_NAMESPACE(cl) {
2020
namespace sycl {
2121
namespace ONEAPI {
22+
namespace experimental {
2223

2324
// Provides functionality to print data from kernels in a C way:
2425
// - On non-host devices this function is directly mapped to printf from
@@ -66,6 +67,7 @@ int printf(const CONSTANT_AS char *__format, Args... args) {
6667
#endif
6768
}
6869

70+
} // namespace experimental
6971
} // namespace ONEAPI
7072
} // namespace sycl
7173
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/include/CL/sycl/ONEAPI/spec_constant.hpp renamed to ‎sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,10 @@
2222

2323
__SYCL_INLINE_NAMESPACE(cl) {
2424
namespace sycl {
25+
class program;
26+
2527
namespace ONEAPI {
28+
namespace experimental {
2629

2730
class spec_const_error : public compile_program_error {
2831
using compile_program_error::compile_program_error;
@@ -56,6 +59,7 @@ template <typename T, typename ID = T> class spec_constant {
5659
}
5760
};
5861

62+
} // namespace experimental
5963
} // namespace ONEAPI
6064
} // namespace sycl
6165
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/include/CL/sycl/program.hpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,13 +8,13 @@
88

99
#pragma once
1010

11+
#include <CL/sycl/ONEAPI/experimental/spec_constant.hpp>
1112
#include <CL/sycl/context.hpp>
1213
#include <CL/sycl/detail/export.hpp>
1314
#include <CL/sycl/detail/kernel_desc.hpp>
1415
#include <CL/sycl/detail/os_util.hpp>
1516
#include <CL/sycl/info/info_desc.hpp>
1617
#include <CL/sycl/kernel.hpp>
17-
#include <CL/sycl/ONEAPI/spec_constant.hpp>
1818
#include <CL/sycl/stl.hpp>
1919

2020
__SYCL_INLINE_NAMESPACE(cl) {
@@ -307,18 +307,18 @@ class __SYCL_EXPORT program {
307307
/// \return a specialization constant instance corresponding to given type ID
308308
/// passed as a template parameter
309309
template <typename ID, typename T>
310-
ONEAPI::spec_constant<T, ID> set_spec_constant(T Cst) {
310+
ONEAPI::experimental::spec_constant<T, ID> set_spec_constant(T Cst) {
311311
constexpr const char *Name = detail::SpecConstantInfo<ID>::getName();
312312
static_assert(std::is_integral<T>::value ||
313313
std::is_floating_point<T>::value,
314314
"unsupported specialization constant type");
315315
#ifdef __SYCL_DEVICE_ONLY__
316316
(void)Cst;
317317
(void)Name;
318-
return ONEAPI::spec_constant<T, ID>();
318+
return ONEAPI::experimental::spec_constant<T, ID>();
319319
#else
320320
set_spec_constant_impl(Name, &Cst, sizeof(T));
321-
return ONEAPI::spec_constant<T, ID>(Cst);
321+
return ONEAPI::experimental::spec_constant<T, ID>(Cst);
322322
#endif // __SYCL_DEVICE_ONLY__
323323
}
324324

‎sycl/source/detail/program_impl.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -484,8 +484,8 @@ vector_class<device> program_impl::get_info<info::program::devices>() const {
484484
void program_impl::set_spec_constant_impl(const char *Name, const void *ValAddr,
485485
size_t ValSize) {
486486
if (MState != program_state::none)
487-
throw cl::sycl::ONEAPI::spec_const_error("Invalid program state",
488-
PI_INVALID_PROGRAM);
487+
throw cl::sycl::ONEAPI::experimental::spec_const_error(
488+
"Invalid program state", PI_INVALID_PROGRAM);
489489
// Reuse cached programs lock as opposed to introducing a new lock.
490490
auto LockGuard = MContext->getKernelProgramCache().acquireCachedPrograms();
491491
spec_constant_impl &SC = SpecConstRegistry[Name];

‎sycl/source/detail/program_manager/program_manager.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9+
#include <CL/sycl/ONEAPI/experimental/spec_constant.hpp>
910
#include <CL/sycl/backend_types.hpp>
1011
#include <CL/sycl/context.hpp>
1112
#include <CL/sycl/detail/common.hpp>
@@ -14,7 +15,6 @@
1415
#include <CL/sycl/detail/util.hpp>
1516
#include <CL/sycl/device.hpp>
1617
#include <CL/sycl/exception.hpp>
17-
#include <CL/sycl/ONEAPI/spec_constant.hpp>
1818
#include <CL/sycl/stl.hpp>
1919
#include <detail/context_impl.hpp>
2020
#include <detail/device_impl.hpp>
@@ -986,7 +986,7 @@ void ProgramManager::flushSpecConstants(const program_impl &Prg,
986986
auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms();
987987
auto It = NativePrograms.find(NativePrg);
988988
if (It == NativePrograms.end())
989-
throw sycl::ONEAPI::spec_const_error(
989+
throw sycl::ONEAPI::experimental::spec_const_error(
990990
"spec constant is set in a program w/o a binary image",
991991
PI_INVALID_OPERATION);
992992
Img = It->second;

‎sycl/test/built-ins/printf.cpp

Lines changed: 19 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -41,7 +41,7 @@ int main() {
4141
Queue.submit([&](handler &CGH) {
4242
CGH.single_task<class integral>([=]() {
4343
// String
44-
ONEAPI::printf(format_hello_world);
44+
ONEAPI::experimental::printf(format_hello_world);
4545
// Due to a bug in Intel CPU Runtime for OpenCL on Windows, information
4646
// printed using such format strings (without %-specifiers) might
4747
// appear in different order if output is redirected to a file or
@@ -50,8 +50,8 @@ int main() {
5050
// CHECK: {{(Hello, World!)?}}
5151

5252
// Integral types
53-
ONEAPI::printf(format_int, (int32_t)123);
54-
ONEAPI::printf(format_int, (int32_t)-123);
53+
ONEAPI::experimental::printf(format_int, (int32_t)123);
54+
ONEAPI::experimental::printf(format_int, (int32_t)-123);
5555
// CHECK: 123
5656
// CHECK-NEXT: -123
5757

@@ -60,8 +60,8 @@ int main() {
6060
// You can declare format string in non-global scope, but in this case
6161
// static keyword is required
6262
static const CONSTANT char format[] = "%f\n";
63-
ONEAPI::printf(format, 33.4f);
64-
ONEAPI::printf(format, -33.4f);
63+
ONEAPI::experimental::printf(format, 33.4f);
64+
ONEAPI::experimental::printf(format, -33.4f);
6565
}
6666
// CHECK-NEXT: 33.4
6767
// CHECK-NEXT: -33.4
@@ -73,23 +73,23 @@ int main() {
7373
using ocl_int4 = cl::sycl::vec<int, 4>::vector_t;
7474
{
7575
static const CONSTANT char format[] = "%v4d\n";
76-
ONEAPI::printf(format, (ocl_int4)v4);
76+
ONEAPI::experimental::printf(format, (ocl_int4)v4);
7777
}
7878

7979
// However, you are still able to print them by-element:
8080
{
81-
ONEAPI::printf(format_vec, (int32_t)v4.w(),
82-
(int32_t)v4.z(), (int32_t)v4.y(),
83-
(int32_t)v4.x());
81+
ONEAPI::experimental::printf(format_vec, (int32_t)v4.w(),
82+
(int32_t)v4.z(), (int32_t)v4.y(),
83+
(int32_t)v4.x());
8484
}
8585
#else
8686
// On host side you always have to print them by-element:
87-
ONEAPI::printf(format_vec, (int32_t)v4.x(),
88-
(int32_t)v4.y(), (int32_t)v4.z(),
89-
(int32_t)v4.w());
90-
ONEAPI::printf(format_vec, (int32_t)v4.w(),
91-
(int32_t)v4.z(), (int32_t)v4.y(),
92-
(int32_t)v4.x());
87+
ONEAPI::experimental::printf(format_vec, (int32_t)v4.x(),
88+
(int32_t)v4.y(), (int32_t)v4.z(),
89+
(int32_t)v4.w());
90+
ONEAPI::experimental::printf(format_vec, (int32_t)v4.w(),
91+
(int32_t)v4.z(), (int32_t)v4.y(),
92+
(int32_t)v4.x());
9393
#endif // __SYCL_DEVICE_ONLY__
9494
// CHECK-NEXT: 5,6,7,8
9595
// CHECK-NEXT: 8,7,6,5
@@ -100,7 +100,7 @@ int main() {
100100
// According to OpenCL spec, argument should be a void pointer
101101
{
102102
static const CONSTANT char format[] = "%p\n";
103-
ONEAPI::printf(format, (void *)Ptr);
103+
ONEAPI::experimental::printf(format, (void *)Ptr);
104104
}
105105
// CHECK-NEXT: {{(0x)?[0-9a-fA-F]+$}}
106106
});
@@ -111,7 +111,7 @@ int main() {
111111
Queue.submit([&](handler &CGH) {
112112
CGH.parallel_for<class stream_string>(range<1>(10), [=](id<1> i) {
113113
// cast to uint64_t to be sure that we pass 64-bit unsigned value
114-
ONEAPI::printf(format_hello_world_2, (uint64_t)i.get(0));
114+
ONEAPI::experimental::printf(format_hello_world_2, (uint64_t)i.get(0));
115115
});
116116
});
117117
Queue.wait();
@@ -127,8 +127,8 @@ int main() {
127127
// CHECK-NEXT: {{[0-9]+}}: Hello, World!
128128
}
129129

130-
// FIXME: strictly check output order once the bug mentioned above is fixed
131-
// CHECK: {{(Hello, World!)?}}
130+
// FIXME: strictly check output order once the bug mentioned above is fixed
131+
// CHECK: {{(Hello, World!)?}}
132132

133133
return 0;
134134
}

‎sycl/test/spec_const/spec_const_hw.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,8 @@ int val = 10;
3535
int get_value() { return val; }
3636

3737
float foo(
38-
const cl::sycl::ONEAPI::spec_constant<float, MyFloatConst> &f32) {
38+
const cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst>
39+
&f32) {
3940
return f32;
4041
}
4142

@@ -66,10 +67,10 @@ int main(int argc, char **argv) {
6667
// TODO make this floating point once supported by the compiler
6768
float goldf = (float)get_value();
6869

69-
cl::sycl::ONEAPI::spec_constant<int32_t, MyInt32Const> i32 =
70+
cl::sycl::ONEAPI::experimental::spec_constant<int32_t, MyInt32Const> i32 =
7071
program1.set_spec_constant<MyInt32Const>(goldi);
7172

72-
cl::sycl::ONEAPI::spec_constant<float, MyFloatConst> f32 =
73+
cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32 =
7374
program2.set_spec_constant<MyFloatConst>(goldf);
7475

7576
program1.build_with_kernel_type<KernelAAAi>();

‎sycl/test/spec_const/spec_const_neg.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -45,7 +45,7 @@ int main(int argc, char **argv) {
4545
<< "\n";
4646
cl::sycl::program program1(q.get_context());
4747

48-
cl::sycl::ONEAPI::spec_constant<int32_t, MyInt32Const> i32 =
48+
cl::sycl::ONEAPI::experimental::spec_constant<int32_t, MyInt32Const> i32 =
4949
program1.set_spec_constant<MyInt32Const>(10);
5050

5151
std::vector<int> veci(1);
@@ -56,7 +56,7 @@ int main(int argc, char **argv) {
5656
try {
5757
// This is an attempt to set a spec constant after the program has been
5858
// built - spec_const_error should be thrown
59-
cl::sycl::ONEAPI::spec_constant<int32_t, MyInt32Const> i32 =
59+
cl::sycl::ONEAPI::experimental::spec_constant<int32_t, MyInt32Const> i32 =
6060
program1.set_spec_constant<MyInt32Const>(10);
6161

6262
cl::sycl::buffer<int, 1> bufi(veci.data(), veci.size());
@@ -69,7 +69,7 @@ int main(int argc, char **argv) {
6969
acci[0] = i32.get();
7070
});
7171
});
72-
} catch (cl::sycl::ONEAPI::spec_const_error &sc_err) {
72+
} catch (cl::sycl::ONEAPI::experimental::spec_const_error &sc_err) {
7373
passed = true;
7474
} catch (cl::sycl::exception &e) {
7575
std::cout << "*** Exception caught: " << e.what() << "\n";

‎sycl/test/spec_const/spec_const_redefine.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -68,9 +68,9 @@ int main(int argc, char **argv) {
6868
for (int i = 0; i < n_sc_sets; i++) {
6969
cl::sycl::program program(q.get_context());
7070
const int *sc_set = &sc_vals[i][0];
71-
cl::sycl::ONEAPI::spec_constant<int32_t, SC0> sc0 =
71+
cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC0> sc0 =
7272
program.set_spec_constant<SC0>(sc_set[0]);
73-
cl::sycl::ONEAPI::spec_constant<int32_t, SC1> sc1 =
73+
cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC1> sc1 =
7474
program.set_spec_constant<SC1>(sc_set[1]);
7575

7676
program.build_with_kernel_type<KernelAAA>();

0 commit comments

Comments
 (0)