Skip to content

Commit 3534bc4

Browse files
mdfazlaymdfazlay
committed
Merge from 'sycl' to 'sycl-web' (6 commits)
 CONFLICT (content): Merge conflict in sycl/test/check_device_code/group_load.cpp CONFLICT (content): Merge conflict in sycl/test/check_device_code/group_store.cpp
2 parents 8219349 + 15929c6 commit 3534bc4

File tree

19 files changed

+403
-297
lines changed

19 files changed

+403
-297
lines changed

‎llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,29 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
3838
auto &DL = M.getDataLayout();
3939
auto &Ctx = M.getContext();
4040

41+
// Fix device global type, by wrapping a structure type
42+
{
43+
assert(KernelMetadata->getValueType()->isArrayTy());
44+
45+
auto *KernelMetadataOld = KernelMetadata;
46+
47+
StructType *StructTypeWithArray = StructType::create(Ctx);
48+
StructTypeWithArray->setBody(KernelMetadataOld->getValueType());
49+
50+
KernelMetadata = new GlobalVariable(
51+
M, StructTypeWithArray, false, GlobalValue::ExternalLinkage,
52+
ConstantStruct::get(StructTypeWithArray,
53+
KernelMetadataOld->getInitializer()),
54+
"", nullptr, GlobalValue::NotThreadLocal, 1); // Global AddressSpace
55+
KernelMetadata->takeName(KernelMetadataOld);
56+
KernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
57+
KernelMetadata->setDSOLocal(true);
58+
KernelMetadata->copyAttributesFrom(KernelMetadataOld);
59+
KernelMetadata->copyMetadata(KernelMetadataOld, 0);
60+
61+
KernelMetadataOld->eraseFromParent();
62+
}
63+
4164
// Fix attributes
4265
KernelMetadata->addAttribute(
4366
"sycl-device-global-size",

‎llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
1818

1919
@__asan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
2020
@__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__asan_kernel to i64), i64 54 }] #2
21-
; CHECK-IR: @__AsanKernelMetadata {{.*}} !spirv.Decorations
21+
; CHECK-IR: @__AsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
2222
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
2323
@__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
2424

‎llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
1818

1919
@__msan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
2020
@__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }] #0
21-
; CHECK-IR: @__MsanKernelMetadata {{.*}} !spirv.Decorations
21+
; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
2222
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
2323
@__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
2424

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc

Lines changed: 21 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -99,6 +99,7 @@ in the group.
9999
* Value type of `InputIteratorT` must be convertible to `OutputT`.
100100
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
101101
and default constructible.
102+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
102103

103104
_Effects_: Loads single element from `in_iter` to `out` by using the `g` group
104105
object to identify memory location as `in_iter` + `g.get_local_linear_id()`.
@@ -129,6 +130,7 @@ in the group.
129130
* Value type of `InputIteratorT` must be convertible to `OutputT`.
130131
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
131132
and default constructible.
133+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
132134

133135
_Effects_: Loads `N` elements from `in_iter` to `out`
134136
using the `g` group object.
@@ -165,6 +167,7 @@ work-group or sub-group.
165167
* Value type of `InputIteratorT` must be convertible to `OutputT`.
166168
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
167169
and default constructible.
170+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
168171

169172
_Effects_: Loads `ElementsPerWorkItem` elements from `in_iter` to `out`
170173
using the `g` group object.
@@ -204,6 +207,7 @@ in the group.
204207
* `InputT` must be convertible to value type of `OutputIteratorT`.
205208
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
206209
and default constructible.
210+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
207211

208212
_Effects_: Stores single element `in` to `out_iter` by using the `g` group
209213
object to identify memory location as `out_iter` + `g.get_local_linear_id()`
@@ -235,6 +239,7 @@ in the group.
235239
* `InputT` must be convertible to value type of `OutputIteratorT`.
236240
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
237241
and default constructible.
242+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
238243

239244
_Effects_: Stores `N` elements from `in` vec to `out_iter`
240245
using the `g` group object.
@@ -273,6 +278,7 @@ work-group or sub-group.
273278
* `InputT` must be convertible to value type of `OutputIteratorT`.
274279
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
275280
and default constructible.
281+
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`
276282

277283
_Effects_: Stores `ElementsPerWorkItem` elements from `in` span to `out_iter`
278284
using the `g` group object.
@@ -370,7 +376,7 @@ Specifies data layout used in group_load/store for `sycl::vec` or fixed-size
370376
arrays functions.
371377

372378
Example:
373-
`group_load(g, input, output_span, data_placement_blocked);`
379+
`group_load(g, input, output_span, properties{data_placement_blocked});`
374380

375381
=== Optimization Properties
376382

@@ -398,7 +404,7 @@ inline constexpr contiguous_memory_key::value_t contiguous_memory;
398404
----
399405

400406
For example, we can assert that `input` is a contiguous iterator:
401-
`group_load(g, input, output_span, contiguous_memory);`
407+
`group_load(g, input, output_span, properties{contiguous_memory});`
402408

403409
If `input` isn't a contiguous iterator, the behavior is undefined.
404410

@@ -432,7 +438,7 @@ inline constexpr full_group_key::value_t full_group;
432438

433439
For example, we can assert that there is no uneven group partition,
434440
so the implementation can rely on `get_max_local_range()` range size:
435-
`group_load(sg, input, output_span, full_group);`
441+
`group_load(sg, input, output_span, properties{full_group});`
436442

437443
If partition is uneven the behavior is undefined.
438444

@@ -466,11 +472,13 @@ q.submit([&](sycl::handler& cgh) {
466472
auto offset = g.get_group_id(0) * g.get_local_range(0) *
467473
items_per_thread;
468474
469-
sycl_exp::group_load(g, input + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
475+
auto props = sycl_exp::properties{sycl_exp::contiguous_memory};
476+
477+
sycl_exp::group_load(g, input + offset, sycl::span{ data }, props);
470478
471479
// Work with data...
472480
473-
sycl_exp::group_store(g, output + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
481+
sycl_exp::group_store(g, output + offset, sycl::span{ data }, props);
474482
});
475483
});
476484
----
@@ -546,11 +554,13 @@ q.submit([&](sycl::handler& cgh) {
546554
sycl_exp::group_with_scratchpad gh{ g,
547555
sycl::span{ buf_ptr, temp_memory_size } };
548556
549-
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
557+
auto props = sycl_exp::properties{sycl_exp::contiguous_memory};
558+
559+
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, props);
550560
551561
// Work with data...
552562
553-
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
563+
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, props);
554564
});
555565
});
556566
----
@@ -583,11 +593,13 @@ q.submit([&](sycl::handler& cgh) {
583593
sycl_exp::group_with_scratchpad gh{ g,
584594
sycl::span{ buf_ptr, temp_memory_size } };
585595
586-
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, sycl_exp::data_placement_striped);
596+
auto striped = sycl_exp::properties{sycl_exp::data_placement_striped};
597+
598+
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, striped);
587599
588600
// Work with data...
589601
590-
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, sycl_exp::data_placement_striped);
602+
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, striped);
591603
});
592604
});
593605
----

‎sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -233,7 +233,8 @@ template <typename Group, typename InputIteratorT, typename OutputT,
233233
std::size_t ElementsPerWorkItem,
234234
typename Properties = decltype(properties())>
235235
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
236-
detail::is_generic_group_v<Group>>
236+
detail::is_generic_group_v<Group> &&
237+
is_property_list_v<Properties>>
237238
group_load(Group g, InputIteratorT in_ptr,
238239
span<OutputT, ElementsPerWorkItem> out, Properties props = {}) {
239240
constexpr bool blocked = detail::isBlocked(props);
@@ -305,7 +306,8 @@ template <typename Group, typename InputT, std::size_t ElementsPerWorkItem,
305306
typename OutputIteratorT,
306307
typename Properties = decltype(properties())>
307308
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
308-
detail::is_generic_group_v<Group>>
309+
detail::is_generic_group_v<Group> &&
310+
is_property_list_v<Properties>>
309311
group_store(Group g, const span<InputT, ElementsPerWorkItem> in,
310312
OutputIteratorT out_ptr, Properties props = {}) {
311313
constexpr bool blocked = detail::isBlocked(props);
@@ -352,7 +354,8 @@ group_store(Group g, const span<InputT, ElementsPerWorkItem> in,
352354
template <typename Group, typename InputIteratorT, typename OutputT,
353355
typename Properties = decltype(properties())>
354356
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
355-
detail::is_generic_group_v<Group>>
357+
detail::is_generic_group_v<Group> &&
358+
is_property_list_v<Properties>>
356359
group_load(Group g, InputIteratorT in_ptr, OutputT &out,
357360
Properties properties = {}) {
358361
group_load(g, in_ptr, span<OutputT, 1>(&out, 1), properties);
@@ -362,7 +365,8 @@ group_load(Group g, InputIteratorT in_ptr, OutputT &out,
362365
template <typename Group, typename InputT, typename OutputIteratorT,
363366
typename Properties = decltype(properties())>
364367
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
365-
detail::is_generic_group_v<Group>>
368+
detail::is_generic_group_v<Group> &&
369+
is_property_list_v<Properties>>
366370
group_store(Group g, const InputT &in, OutputIteratorT out_ptr,
367371
Properties properties = {}) {
368372
group_store(g, span<const InputT, 1>(&in, 1), out_ptr, properties);
@@ -372,7 +376,8 @@ group_store(Group g, const InputT &in, OutputIteratorT out_ptr,
372376
template <typename Group, typename InputIteratorT, typename OutputT, int N,
373377
typename Properties = decltype(properties())>
374378
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
375-
detail::is_generic_group_v<Group>>
379+
detail::is_generic_group_v<Group> &&
380+
is_property_list_v<Properties>>
376381
group_load(Group g, InputIteratorT in_ptr, sycl::vec<OutputT, N> &out,
377382
Properties properties = {}) {
378383
group_load(g, in_ptr, span<OutputT, N>(&out[0], N), properties);
@@ -382,7 +387,8 @@ group_load(Group g, InputIteratorT in_ptr, sycl::vec<OutputT, N> &out,
382387
template <typename Group, typename InputT, int N, typename OutputIteratorT,
383388
typename Properties = decltype(properties())>
384389
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
385-
detail::is_generic_group_v<Group>>
390+
detail::is_generic_group_v<Group> &&
391+
is_property_list_v<Properties>>
386392
group_store(Group g, const sycl::vec<InputT, N> &in, OutputIteratorT out_ptr,
387393
Properties properties = {}) {
388394
group_store(g, span<const InputT, N>(&in[0], N), out_ptr, properties);

‎sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
// REQUIRES: aspect-ext_intel_matrix
1111
// REQUIRES-INTEL-DRIVER: lin: 30049, win: 101.4943
1212

13-
// XFAIL: !igc-dev
13+
// XFAIL: !igc-dev && run-mode
1414
// XFAIL-TRACKER: GSD-6376
1515

1616
// RUN: %{build} -o %t.out
@@ -20,4 +20,4 @@
2020

2121
#define SG_SZ 32
2222

23-
#include "../get_coordinate_ops_impl.hpp"
23+
#include "../get_coordinate_ops_impl.hpp"

‎sycl/test-e2e/MemorySanitizer/check_buffer.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,9 +4,6 @@
44
// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
55
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
66

7-
// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
8-
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
9-
107
#include <sycl/detail/core.hpp>
118

129
__attribute__((noinline)) long long foo(int data1, long long data2) {

‎sycl/test-e2e/MemorySanitizer/check_call.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,9 +4,6 @@
44
// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
55
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
66

7-
// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
8-
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
9-
107
#include <sycl/detail/core.hpp>
118
#include <sycl/usm.hpp>
129

‎sycl/test-e2e/MemorySanitizer/check_divide.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4,9 +4,6 @@
44
// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
55
// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
66

7-
// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
8-
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
9-
107
#include <sycl/detail/core.hpp>
118
#include <sycl/usm.hpp>
129

‎sycl/test-e2e/MemorySanitizer/lit.local.cfg

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,8 @@
1+
# TRACKER: https://github.com/intel/llvm/issues/16184
2+
has_arch_gpu_intel_pvc = any('arch-intel_gpu_pvc' in T for T in config.sycl_dev_features.values())
3+
if not has_arch_gpu_intel_pvc:
4+
config.unsupported_features += ['gpu']
5+
16
config.substitutions.append(
27
("%device_msan_flags", "-Xarch_device -fsanitize=memory")
38
)

‎sycl/test-e2e/OnlineCompiler/online_compiler_OpenCL.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// REQUIRES: opencl, opencl_icd, cm-compiler
2-
// XFAIL: gpu || cpu
2+
// XFAIL: gpu || cpu || accelerator
33
// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16406
44
// RUN: %{build} -Wno-error=deprecated-declarations -DRUN_KERNELS %opencl_lib -o %t.out
55
// RUN: %{run} %t.out

‎sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
// REQUIRES: build-and-run-mode
44

55
// DEFINE: %{link-flags}=%if cl_options %{ /clang:-ld3d12 /clang:-ldxgi /clang:-ldxguid %} %else %{ -ld3d12 -ldxgi -ldxguid %}
6-
// RUN: %{build} %{link-flags} -o %t.out
6+
// RUN: %{build} %{link-flags} -o %t.out %if any-device-is-level_zero %{ -DDISABLE_UNORM_TESTS %}
77
// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
88

99
#pragma clang diagnostic ignored "-Waddress-of-temporary"
@@ -733,8 +733,10 @@ int main() {
733733
validated &=
734734
runTest<1, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
735735
globalSize1, localSize1);
736+
#ifndef DISABLE_UNORM_TESTS
736737
validated &= runTest<1, uint8_t, 4>(
737738
device, sycl::image_channel_type::unorm_int8, globalSize1, localSize1);
739+
#endif
738740
validated &= runTest<1, float, 1>(device, sycl::image_channel_type::fp32,
739741
globalSize1, localSize1);
740742
validated &= runTest<1, sycl::half, 2>(device, sycl::image_channel_type::fp16,
@@ -752,8 +754,10 @@ int main() {
752754
validated &=
753755
runTest<2, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
754756
globalSize2[0], {16, 16});
757+
#ifndef DISABLE_UNORM_TESTS
755758
validated &= runTest<2, uint8_t, 4>(
756759
device, sycl::image_channel_type::unorm_int8, globalSize2[1], {16, 8});
760+
#endif
757761
validated &= runTest<2, float, 1>(device, sycl::image_channel_type::fp32,
758762
globalSize2[2], {16, 8});
759763
validated &= runTest<2, sycl::half, 2>(device, sycl::image_channel_type::fp16,
@@ -774,8 +778,10 @@ int main() {
774778
validated &=
775779
runTest<3, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
776780
globalSize3[0], {16, 16, 1});
781+
#ifndef DISABLE_UNORM_TESTS
777782
validated &= runTest<3, uint8_t, 4>(
778783
device, sycl::image_channel_type::unorm_int8, globalSize3[1], {16, 8, 2});
784+
#endif
779785
validated &= runTest<3, float, 1>(device, sycl::image_channel_type::fp32,
780786
globalSize3[2], {16, 8, 1});
781787
validated &= runTest<3, sycl::half, 2>(device, sycl::image_channel_type::fp16,

‎sycl/test-e2e/bindless_images/helpers/common.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,8 @@ constexpr sycl::vec<DType, NChannel> init_vector(DType val) {
8282
return sycl::vec<DType, NChannel>{val};
8383
} else if constexpr (NChannel == 2) {
8484
return sycl::vec<DType, NChannel>{val, val};
85+
} else if constexpr (NChannel == 3) {
86+
return sycl::vec<DType, NChannel>{val, val, val};
8587
} else if constexpr (NChannel == 4) {
8688
return sycl::vec<DType, NChannel>{val, val, val, val};
8789
} else {

‎sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -483,6 +483,12 @@ bool run_tests() {
483483
valid &= run_test<2, sycl::half, 2, sycl::image_channel_type::fp16,
484484
sycl::image_channel_order::rg, class fp16_2d_c2>(
485485
{1920, 1080}, {16, 8}, 0);
486+
valid &= run_test<2, sycl::half, 3, sycl::image_channel_type::fp16,
487+
sycl::image_channel_order::rgb, class fp16_2d_c3>(
488+
{2048, 2048}, {16, 16}, 0);
489+
valid &= run_test<2, uint8_t, 3, sycl::image_channel_type::unorm_int8,
490+
sycl::image_channel_order::rgb, class unorm_int8_2d_c3>(
491+
{2048, 2048}, {16, 16}, 0);
486492
valid &= run_test<2, sycl::half, 4, sycl::image_channel_type::fp16,
487493
sycl::image_channel_order::rgba, class fp16_2d_c4>(
488494
{2048, 2048}, {16, 16}, 0);

‎sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp

Lines changed: 5 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -590,11 +590,6 @@ bool run_all() {
590590
sycl::image_channel_order::rgba, class fp16_3d_c4>(
591591
{2048, 2048, 4}, {16, 16, 1}, seed);
592592

593-
printString("Running 3D unorm_int8_c4\n");
594-
valid &= run_test<3, uint8_t, 4, sycl::image_channel_type::unorm_int8,
595-
sycl::image_channel_order::rgba, class unorm_int8_3d_c4>(
596-
{2048, 2048, 2}, {16, 16, 1}, seed);
597-
598593
printString("Running 2D float\n");
599594
valid &= run_test<2, float, 1, sycl::image_channel_type::fp32,
600595
sycl::image_channel_order::r, class fp32_2d_c1>(
@@ -608,10 +603,12 @@ bool run_all() {
608603
sycl::image_channel_order::rgba, class fp16_2d_c4>(
609604
{2048, 2048}, {16, 16}, seed);
610605

611-
printString("Running 2D unorm_int8_c4\n");
612-
valid &= run_test<2, uint8_t, 4, sycl::image_channel_type::unorm_int8,
613-
sycl::image_channel_order::rgba, class unorm_int8_2d_c4>(
606+
// 3-channels
607+
printString("Running 2D half3\n");
608+
valid &= run_test<2, sycl::half, 3, sycl::image_channel_type::fp16,
609+
sycl::image_channel_order::rgb, class fp16_2d_c3>(
614610
{2048, 2048}, {2, 2}, seed);
611+
615612
#else
616613
printString("Running 3D uint4\n");
617614
valid &= run_test<3, uint32_t, 4, sycl::image_channel_type::signed_int32,

0 commit comments

Comments
 (0)