From d4a3c995a1f855ee186852f35fa379cb8e7d5eb0 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 27 Sep 2021 17:06:50 +0300 Subject: [PATCH 1/5] Add test for vec with std::byte --- SYCL/Basic/vector_byte.cpp | 296 +++++++++++++++++++++++++++++++++++++ 1 file changed, 296 insertions(+) create mode 100644 SYCL/Basic/vector_byte.cpp diff --git a/SYCL/Basic/vector_byte.cpp b/SYCL/Basic/vector_byte.cpp new file mode 100644 index 0000000000..ef93f68ab4 --- /dev/null +++ b/SYCL/Basic/vector_byte.cpp @@ -0,0 +1,296 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +//==---------- vector_byte.cpp - SYCL vec<> for std::byte test -------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#define SYCL_SIMPLE_SWIZZLES +#include + +int main() { + std::byte bt{7}; + // constructors + sycl::vec vb1 (bt); + sycl::vec vb2 {bt, bt}; + sycl::vec vb3 {bt, bt, bt}; + sycl::vec vb4 {bt, bt, bt, bt}; + sycl::vec vb8 {bt, bt, bt, bt, bt, bt, bt, bt}; + sycl::vec vb16 {bt, bt, bt, std::byte{2}, bt, bt, bt, bt, bt, + bt, bt, bt, bt, bt, bt, bt}; + + { + // operator[] + assert(vb16[3] == std::byte{2}); + // explicit conversion + std::byte(vb1.x()); + std::byte b = vb1; + + // operator= + auto vb4op = vb4; + vb1 = std::byte{3}; + } + + // convert() and as() + { + sycl::vec vi2(1,1); + auto cnv = vi2.convert(); + auto cnv2 = vb1.convert(); + + auto asint = vb2.template as>(); + auto asbyte = vi2.template as>(); + } + + // load() and store() + { + std::vector std_vec(8, bt); + sycl::buffer Buf(std_vec.data(), sycl::range<1>(8)); + + sycl::queue Queue; + Queue.submit([&](sycl::handler &cgh) { + auto Acc = Buf.get_access(cgh); + cgh.single_task([=]() { + // load + sycl::multi_ptr mp( + &Acc[0]); + sycl::vec sycl_vec; + sycl_vec.load(0, mp); + std::byte bc{7}; + assert(sycl_vec[0] == bc); + assert(sycl_vec[1] == bc); + sycl_vec[0] = std::byte{2}; + + //store + sycl_vec.store(0, mp); + assert(Acc[0] == std::byte{2}); + assert(Acc[1] == bc); + }); + }); + } + + // swizzle + { + auto swizzled_vec = vb8.lo(); + + auto sw = vb8.template swizzle().template as >(); + auto swbyte = sw.template as>(); + auto swbyte2 = swizzled_vec.template as >(); + + // hi/lo, even/odd + sycl::vec vbsw(std::byte{0}, std::byte{1}, std::byte{2}, std::byte{3}); + sycl::vec vbswhi = vbsw.hi(); + assert (vbswhi[0] == std::byte{2}); + vbswhi = vbsw.lo(); + vbswhi = vbsw.odd(); + vbswhi = vbsw.even(); + } + + //operatorOP for vec and for swizzle + { + sycl::vec vop1 {std::byte{4},std::byte{9},std::byte{25}}; + sycl::vec vop2 {std::byte{2},std::byte{3},std::byte{5}}; + sycl::vec vop3 {std::byte{5},std::byte{6},std::byte{2}, std::byte{3}}; + + // binary op for 2 vec + auto vop = vop1 + vop2; + assert (vop[0] == std::byte{6}); + vop = vop1 - vop2; + vop = vop1 * vop2; + vop = vop1 / vop2; + assert (vop[0] == std::byte{2}); + vop = vop1 % vop2; + + // binary op for 2 swizzle + auto swlo = vop3.lo(); + auto swhi = vop3.hi(); + auto swplus = swlo + swhi; + sycl::vec vec_test = swplus; + assert(vec_test.x() == std::byte{7} && vec_test.y() == std::byte{9}); + auto swominus = swlo - swhi; + auto swmul = swlo * swhi; + vec_test = swmul; + assert(vec_test.x() == std::byte{10} && vec_test.y() == std::byte{18}); + auto swdiv = swlo / swhi; + + // binary op for 1 vec + vop = vop1 + std::byte{3}; + vop = vop1 - std::byte{3}; + assert (vop[1] == std::byte{6}); + vop = vop1 * std::byte{3}; + vop = vop1 / std::byte{3}; + vop = vop1 % std::byte{3}; + assert (vop[0] == std::byte{1}); + + vop = std::byte{3} + vop1; + assert (vop[0] == std::byte{7}); + vop = std::byte{3} - vop1; + vop = std::byte{3} * vop1; + assert (vop[2] == std::byte{75}); + vop = std::byte{3} / vop1; + + // binary op for 1 swizzle + auto swplus1 = swlo + std::byte{3}; + auto swminus1 = swlo - std::byte{3}; + vec_test = swminus1; + assert(vec_test.x() == std::byte{2} && vec_test.y() == std::byte{3}); + auto swmul1 = swlo * std::byte{3}; + auto swdiv1 = swlo / std::byte{3}; + vec_test = swdiv1; + assert(vec_test.x() == std::byte{1} && vec_test.y() == std::byte{2}); + + auto swplus2 = std::byte{3} + swlo; + vec_test = swplus2; + assert(vec_test.x() == std::byte{8} && vec_test.y() == std::byte{9}); + auto swminus2 = std::byte{3} - swlo; + auto swmul2 = std::byte{3} * swlo; + vec_test = swmul2; + assert(vec_test.x() == std::byte{15} && vec_test.y() == std::byte{18}); + auto swdiv2 = std::byte{3} / swlo; + + // operatorOP= for 2 vec + sycl::vec vbuf {std::byte{4},std::byte{5},std::byte{6}}; + vop = vbuf += vop1; + assert (vop[0] == std::byte{8}); + vop = vbuf -= vop1; + vop = vbuf *= vop1; + vop = vbuf /= vop1; + vop = vbuf %= vop1; + + // operatorOP= for 2 swizzle + swlo += swhi; + swlo -= swhi; + vec_test = swlo; + assert(vec_test.x() == std::byte{5} && vec_test.y() == std::byte{6}); + swlo *= swhi; + swlo /= swhi; + swlo %= swhi; + + // operatorOP= for 1 vec + vop = vop1 += std::byte{3}; + assert(vop[0] == std::byte{7}); + vop = vop1 -= std::byte{3}; + vop = vop1 *= std::byte{3}; + vop = vop1 /= std::byte{3}; + vop = vop1 %= std::byte{3}; + + // operatorOP= for 1 swizzle + + swlo += std::byte{3}; + swlo -= std::byte{1}; + vec_test = swlo; + assert(vec_test.x() == std::byte{3} && vec_test.y() == std::byte{2}); + swlo *= std::byte{3}; + swlo /= std::byte{3}; + swlo %= std::byte{3}; + + // unary operator++ and -- for vec + vop1 = sycl::vec(std::byte{4},std::byte{9},std::byte{25}); + vop1++; + vop1--; + vop = ++vop1; + assert(vop[2] == std::byte{26}); + --vop1; + + // unary operator++ and -- for swizzle + swlo++; + swlo--; + vec_test = swlo; + assert(vec_test.x() == std::byte{0} && vec_test.y() == std::byte{2}); + + // logical binary op for 2 vec + vop = vop1 & vop2; + vop = vop1 | vop2; + vop = vop1 ^ vop2; + + // logical binary op for 2 swizzle + auto swand = swlo & swhi; + auto swor = swlo | swhi; + auto swxor = swlo ^ swhi; + + // logical binary op for 1 vec + vop = vop1 & std::byte{3}; + vop = vop1 | std::byte{3}; + vop = vop1 ^ std::byte{3}; + vop = std::byte{3} & vop1; + vop = std::byte{3} | vop1; + vop = std::byte{3} ^ vop1; + + // logical binary op for 1 swizzle + auto swand2 = swlo & std::byte{3}; + auto swor2 = swlo | std::byte{3}; + auto swxor2 = swlo ^ std::byte{3}; + + auto swand3 = std::byte{3} & swlo; + auto swor3 = std::byte{3} | swlo; + auto swxor3 = std::byte{3} ^ swlo; + + // bit binary op for 2 vec + vop = vop1 && vop2; + vop = vop1 || vop2; + vop = vop1 >> vop2; + vop = vop1 << vop2; + + vop = vop1 >> std::byte{3}; + vop = vop1 << std::byte{3}; + vop = std::byte{3} >> vop1; + vop = std::byte{3} << vop1; + + // bit binary op for 2 swizzle + swlo >> swhi; + swlo << swhi; + swlo >> std::byte{3}; + swlo << std::byte{3}; + auto right = std::byte{3} >> swhi; + auto left = std::byte{3} << swhi; + + // condition op for 2 vec + auto vres = vop1 == vop2; + vres = vop1 != vop2; + vres = vop1 > vop2; + vres = vop1 < vop2; + vres = vop1 >= vop2; + vres = vop1 <= vop2; + + vres = vop1 == std::byte{3}; + vres = vop1 != std::byte{3}; + vres = vop1 > std::byte{3}; + vres = vop1 < std::byte{3}; + vres = vop1 >= std::byte{3}; + vres = vop1 <= std::byte{3}; + + vres = std::byte{3} == vop1; + vres = std::byte{3} != vop1; + vres = std::byte{3} > vop1; + vres = std::byte{3} < vop1; + vres = std::byte{3} >= vop1; + vres = std::byte{3} <= vop1; + + // condition op for 2 swizzle + auto swres = swhi == swlo; + auto swres1 = swhi != swlo; + auto swres2 = swhi > swlo; + auto swres3 = swhi < swlo; + auto swres4 = swhi >= swlo; + auto swres5 = swhi <= swlo; + auto swres6 = swhi == std::byte{3}; + auto swres7 = swhi != std::byte{3}; + auto swres8 = swhi > std::byte{3}; + auto swres9 = swhi < std::byte{3}; + auto swres10 = swhi >= std::byte{3}; + auto swres11 = swhi <= std::byte{3}; + + sycl::vec voptest {std::byte{4},std::byte{9},std::byte{25}}; + auto bitv1 = ~vop3; + auto bitv2 = !vop3; + auto bitw = ~swhi; + } + + return 0; +} \ No newline at end of file From 935c07a4e97a022ce8dc2e0134600fe7a9c6f0e8 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 27 Sep 2021 17:28:15 +0300 Subject: [PATCH 2/5] Clang-format fix --- SYCL/Basic/vector_byte.cpp | 71 ++++++++++++++++++++------------------ 1 file changed, 37 insertions(+), 34 deletions(-) diff --git a/SYCL/Basic/vector_byte.cpp b/SYCL/Basic/vector_byte.cpp index ef93f68ab4..90a522d817 100644 --- a/SYCL/Basic/vector_byte.cpp +++ b/SYCL/Basic/vector_byte.cpp @@ -18,13 +18,13 @@ int main() { std::byte bt{7}; // constructors - sycl::vec vb1 (bt); - sycl::vec vb2 {bt, bt}; - sycl::vec vb3 {bt, bt, bt}; - sycl::vec vb4 {bt, bt, bt, bt}; - sycl::vec vb8 {bt, bt, bt, bt, bt, bt, bt, bt}; - sycl::vec vb16 {bt, bt, bt, std::byte{2}, bt, bt, bt, bt, bt, - bt, bt, bt, bt, bt, bt, bt}; + sycl::vec vb1(bt); + sycl::vec vb2{bt, bt}; + sycl::vec vb3{bt, bt, bt}; + sycl::vec vb4{bt, bt, bt, bt}; + sycl::vec vb8{bt, bt, bt, bt, bt, bt, bt, bt}; + sycl::vec vb16{bt, bt, bt, std::byte{2}, bt, bt, bt, bt, + bt, bt, bt, bt, bt, bt, bt, bt}; { // operator[] @@ -40,7 +40,7 @@ int main() { // convert() and as() { - sycl::vec vi2(1,1); + sycl::vec vi2(1, 1); auto cnv = vi2.convert(); auto cnv2 = vb1.convert(); @@ -52,14 +52,14 @@ int main() { { std::vector std_vec(8, bt); sycl::buffer Buf(std_vec.data(), sycl::range<1>(8)); - + sycl::queue Queue; Queue.submit([&](sycl::handler &cgh) { auto Acc = Buf.get_access(cgh); cgh.single_task([=]() { // load - sycl::multi_ptr mp( - &Acc[0]); + sycl::multi_ptr + mp(&Acc[0]); sycl::vec sycl_vec; sycl_vec.load(0, mp); std::byte bc{7}; @@ -67,7 +67,7 @@ int main() { assert(sycl_vec[1] == bc); sycl_vec[0] = std::byte{2}; - //store + // store sycl_vec.store(0, mp); assert(Acc[0] == std::byte{2}); assert(Acc[1] == bc); @@ -79,32 +79,35 @@ int main() { { auto swizzled_vec = vb8.lo(); - auto sw = vb8.template swizzle().template as >(); + auto sw = vb8.template swizzle() + .template as>(); auto swbyte = sw.template as>(); - auto swbyte2 = swizzled_vec.template as >(); + auto swbyte2 = swizzled_vec.template as>(); // hi/lo, even/odd - sycl::vec vbsw(std::byte{0}, std::byte{1}, std::byte{2}, std::byte{3}); + sycl::vec vbsw(std::byte{0}, std::byte{1}, std::byte{2}, + std::byte{3}); sycl::vec vbswhi = vbsw.hi(); - assert (vbswhi[0] == std::byte{2}); + assert(vbswhi[0] == std::byte{2}); vbswhi = vbsw.lo(); vbswhi = vbsw.odd(); vbswhi = vbsw.even(); } - //operatorOP for vec and for swizzle + // operatorOP for vec and for swizzle { - sycl::vec vop1 {std::byte{4},std::byte{9},std::byte{25}}; - sycl::vec vop2 {std::byte{2},std::byte{3},std::byte{5}}; - sycl::vec vop3 {std::byte{5},std::byte{6},std::byte{2}, std::byte{3}}; + sycl::vec vop1{std::byte{4}, std::byte{9}, std::byte{25}}; + sycl::vec vop2{std::byte{2}, std::byte{3}, std::byte{5}}; + sycl::vec vop3{std::byte{5}, std::byte{6}, std::byte{2}, + std::byte{3}}; // binary op for 2 vec auto vop = vop1 + vop2; - assert (vop[0] == std::byte{6}); + assert(vop[0] == std::byte{6}); vop = vop1 - vop2; vop = vop1 * vop2; vop = vop1 / vop2; - assert (vop[0] == std::byte{2}); + assert(vop[0] == std::byte{2}); vop = vop1 % vop2; // binary op for 2 swizzle @@ -122,17 +125,17 @@ int main() { // binary op for 1 vec vop = vop1 + std::byte{3}; vop = vop1 - std::byte{3}; - assert (vop[1] == std::byte{6}); + assert(vop[1] == std::byte{6}); vop = vop1 * std::byte{3}; vop = vop1 / std::byte{3}; vop = vop1 % std::byte{3}; - assert (vop[0] == std::byte{1}); + assert(vop[0] == std::byte{1}); vop = std::byte{3} + vop1; - assert (vop[0] == std::byte{7}); + assert(vop[0] == std::byte{7}); vop = std::byte{3} - vop1; vop = std::byte{3} * vop1; - assert (vop[2] == std::byte{75}); + assert(vop[2] == std::byte{75}); vop = std::byte{3} / vop1; // binary op for 1 swizzle @@ -155,9 +158,9 @@ int main() { auto swdiv2 = std::byte{3} / swlo; // operatorOP= for 2 vec - sycl::vec vbuf {std::byte{4},std::byte{5},std::byte{6}}; + sycl::vec vbuf{std::byte{4}, std::byte{5}, std::byte{6}}; vop = vbuf += vop1; - assert (vop[0] == std::byte{8}); + assert(vop[0] == std::byte{8}); vop = vbuf -= vop1; vop = vbuf *= vop1; vop = vbuf /= vop1; @@ -172,7 +175,7 @@ int main() { swlo /= swhi; swlo %= swhi; - // operatorOP= for 1 vec + // operatorOP= for 1 vec vop = vop1 += std::byte{3}; assert(vop[0] == std::byte{7}); vop = vop1 -= std::byte{3}; @@ -191,7 +194,7 @@ int main() { swlo %= std::byte{3}; // unary operator++ and -- for vec - vop1 = sycl::vec(std::byte{4},std::byte{9},std::byte{25}); + vop1 = sycl::vec(std::byte{4}, std::byte{9}, std::byte{25}); vop1++; vop1--; vop = ++vop1; @@ -215,10 +218,10 @@ int main() { auto swxor = swlo ^ swhi; // logical binary op for 1 vec - vop = vop1 & std::byte{3}; + vop = vop1 & std::byte{3}; vop = vop1 | std::byte{3}; vop = vop1 ^ std::byte{3}; - vop = std::byte{3} & vop1; + vop = std::byte{3} & vop1; vop = std::byte{3} | vop1; vop = std::byte{3} ^ vop1; @@ -236,7 +239,7 @@ int main() { vop = vop1 || vop2; vop = vop1 >> vop2; vop = vop1 << vop2; - + vop = vop1 >> std::byte{3}; vop = vop1 << std::byte{3}; vop = std::byte{3} >> vop1; @@ -286,7 +289,7 @@ int main() { auto swres10 = swhi >= std::byte{3}; auto swres11 = swhi <= std::byte{3}; - sycl::vec voptest {std::byte{4},std::byte{9},std::byte{25}}; + sycl::vec voptest{std::byte{4}, std::byte{9}, std::byte{25}}; auto bitv1 = ~vop3; auto bitv2 = !vop3; auto bitw = ~swhi; From 9eb50ce447411a30eb104cff6281a608cdfa5867 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 1 Oct 2021 15:20:10 +0300 Subject: [PATCH 3/5] Added c++17 compile options --- SYCL/Basic/vector_byte.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Basic/vector_byte.cpp b/SYCL/Basic/vector_byte.cpp index 90a522d817..5d47faa2ce 100644 --- a/SYCL/Basic/vector_byte.cpp +++ b/SYCL/Basic/vector_byte.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %clangxx -fsycl -std=c++17 -fsycl-targets=%sycl_triple %s -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From b62aaefdba62d12a72875f839906d4894f35e82c Mon Sep 17 00:00:00 2001 From: mdimakov Date: Tue, 5 Oct 2021 12:07:05 +0300 Subject: [PATCH 4/5] Remove assert from kernel --- SYCL/Basic/vector_byte.cpp | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/SYCL/Basic/vector_byte.cpp b/SYCL/Basic/vector_byte.cpp index 5d47faa2ce..9057d8cb13 100644 --- a/SYCL/Basic/vector_byte.cpp +++ b/SYCL/Basic/vector_byte.cpp @@ -49,8 +49,8 @@ int main() { } // load() and store() + std::vector std_vec(8, bt); { - std::vector std_vec(8, bt); sycl::buffer Buf(std_vec.data(), sycl::range<1>(8)); sycl::queue Queue; @@ -62,18 +62,17 @@ int main() { mp(&Acc[0]); sycl::vec sycl_vec; sycl_vec.load(0, mp); - std::byte bc{7}; - assert(sycl_vec[0] == bc); - assert(sycl_vec[1] == bc); sycl_vec[0] = std::byte{2}; + Acc[1] = std::byte{10}; // store sycl_vec.store(0, mp); - assert(Acc[0] == std::byte{2}); - assert(Acc[1] == bc); + }); - }); + }).wait(); } + assert(std_vec[0] == std::byte{2}); + assert(std_vec[1] == std::byte{7}); // swizzle { From e0195640fc9409be233e08533e489f0d0d96e532 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Tue, 5 Oct 2021 12:14:50 +0300 Subject: [PATCH 5/5] Clang-format fix --- SYCL/Basic/vector_byte.cpp | 34 ++++++++++++++++++---------------- 1 file changed, 18 insertions(+), 16 deletions(-) diff --git a/SYCL/Basic/vector_byte.cpp b/SYCL/Basic/vector_byte.cpp index 9057d8cb13..b559487149 100644 --- a/SYCL/Basic/vector_byte.cpp +++ b/SYCL/Basic/vector_byte.cpp @@ -54,22 +54,24 @@ int main() { sycl::buffer Buf(std_vec.data(), sycl::range<1>(8)); sycl::queue Queue; - Queue.submit([&](sycl::handler &cgh) { - auto Acc = Buf.get_access(cgh); - cgh.single_task([=]() { - // load - sycl::multi_ptr - mp(&Acc[0]); - sycl::vec sycl_vec; - sycl_vec.load(0, mp); - sycl_vec[0] = std::byte{2}; - Acc[1] = std::byte{10}; - - // store - sycl_vec.store(0, mp); - - }); - }).wait(); + Queue + .submit([&](sycl::handler &cgh) { + auto Acc = Buf.get_access(cgh); + cgh.single_task([=]() { + // load + sycl::multi_ptr + mp(&Acc[0]); + sycl::vec sycl_vec; + sycl_vec.load(0, mp); + sycl_vec[0] = std::byte{2}; + Acc[1] = std::byte{10}; + + // store + sycl_vec.store(0, mp); + }); + }) + .wait(); } assert(std_vec[0] == std::byte{2}); assert(std_vec[1] == std::byte{7});