From 006e43d76bb785771700e15765521dcb9f07ab41 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Tue, 14 Mar 2023 06:17:22 -0700 Subject: [PATCH 1/4] Adds tests for atomic_fence_capabilities. Signed-off-by: Maronas, Marcos --- SYCL/AtomicRef/atomic_fence_capabilities.cpp | 109 +++++++++++++++++++ SYCL/AtomicRef/atomic_fence_capabilities.h | 18 +++ 2 files changed, 127 insertions(+) create mode 100644 SYCL/AtomicRef/atomic_fence_capabilities.cpp create mode 100644 SYCL/AtomicRef/atomic_fence_capabilities.h diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.cpp b/SYCL/AtomicRef/atomic_fence_capabilities.cpp new file mode 100644 index 0000000000..16b8cc7bdc --- /dev/null +++ b/SYCL/AtomicRef/atomic_fence_capabilities.cpp @@ -0,0 +1,109 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +// NOTE: General tests for atomic fence capabilities. + +#include "atomic_fence_capabilities.h" +#include +#include +using namespace sycl; + +void checkFenceBehaviour(memory_order order, memory_scope scope) { + auto q = queue(); + memory_order order_write = order; + memory_order order_read = order; + if (memory_order::release == order_read) { + order_read = memory_order::acquire; + } + + // Count of retries in the check cycle + constexpr size_t RETRY_COUNT = 256; + constexpr int expected_val = 42; + + bool res = true; + int sync = 0; + int data = 0; + int value = expected_val; + + // These global_range and local_range values provide a check in one group + // when test_type = single_group, and between four groups when + // test_type = between_groups + range<1> global_range(2); + range<1> local_range(2); + + { + buffer res_buf(&res, range<1>(1)); + buffer sync_buffer(&sync, range<1>(1)); + buffer data_buffer(&data, range<1>(1)); + q.submit([&](handler &cgh) { + auto res_acc = res_buf.template get_access(cgh); + auto sync_flag_acc = + sync_buffer.template get_access(cgh); + auto data_acc = + data_buffer.template get_access(cgh); + cgh.parallel_for( + nd_range<1>(global_range, local_range), [=](nd_item<1> nditem) { + atomic_ref + sync_flag(sync_flag_acc[0]); + int *data = &data_acc[0]; + // Only one nditem should perform non-atomic write. + // All other nditems should perform non-atomic + // reads + if (nditem.get_global_linear_id() == 0) { + // Non-atomic write to data + *data = value; + // Used atomic_fence to guarantee the order + // instructions execution + atomic_fence(order_write, scope); + // Used atomic sync flag to avoid data racing + sync_flag = 1; + } else { + bool write_happened = false; + for (size_t i = 0; i < RETRY_COUNT; i++) { + if (sync_flag == 1) { + write_happened = true; + break; + } + } + atomic_fence(order_read, scope); + // After the fence safe non-atomic reading + if (write_happened) { + // Non-atomic read of data + if (*data != value) + res_acc[0] = false; + } + } + }); + }); + } + assert(res); +} + +int main() { + queue q; + + std::vector supported_memory_orders = + q.get_device().get_info(); + + // Relaxed, acquire, release and acq_rel memory order must be supported. + assert(is_supported_order(supported_memory_orders, memory_order::relaxed)); + assert(is_supported_order(supported_memory_orders, memory_order::acquire)); + assert(is_supported_order(supported_memory_orders, memory_order::release)); + assert(is_supported_order(supported_memory_orders, memory_order::acq_rel)); + + std::vector supported_memory_scopes = + q.get_device().get_info(); + + // Work_group, sub_group and work_item memory order must be supported. + assert(is_supported_scope(supported_memory_scopes, memory_scope::work_item)); + assert(is_supported_scope(supported_memory_scopes, memory_scope::sub_group)); + assert(is_supported_scope(supported_memory_scopes, memory_scope::work_group)); + + for (auto order : supported_memory_orders) + for (auto scope : supported_memory_scopes) + checkFenceBehaviour(order, scope); + + std::cout << "Test passed." << std::endl; +} diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.h b/SYCL/AtomicRef/atomic_fence_capabilities.h new file mode 100644 index 0000000000..75c8eb9255 --- /dev/null +++ b/SYCL/AtomicRef/atomic_fence_capabilities.h @@ -0,0 +1,18 @@ +#pragma once + +#include +#include + +using namespace sycl; + +bool is_supported_order(const std::vector &capabilities, + memory_order mem_order) { + return std::find(capabilities.begin(), capabilities.end(), mem_order) != + capabilities.end(); +} + +bool is_supported_scope(const std::vector &capabilities, + memory_scope mem_scope) { + return std::find(capabilities.begin(), capabilities.end(), mem_scope) != + capabilities.end(); +} From 9196a965794d714e022b73ef7337565ba6300070 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Tue, 14 Mar 2023 08:06:37 -0700 Subject: [PATCH 2/4] Fixes clang-format issue. Signed-off-by: Maronas, Marcos --- SYCL/AtomicRef/atomic_fence_capabilities.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.h b/SYCL/AtomicRef/atomic_fence_capabilities.h index 75c8eb9255..6688ecc320 100644 --- a/SYCL/AtomicRef/atomic_fence_capabilities.h +++ b/SYCL/AtomicRef/atomic_fence_capabilities.h @@ -6,13 +6,13 @@ using namespace sycl; bool is_supported_order(const std::vector &capabilities, - memory_order mem_order) { + memory_order mem_order) { return std::find(capabilities.begin(), capabilities.end(), mem_order) != capabilities.end(); } bool is_supported_scope(const std::vector &capabilities, - memory_scope mem_scope) { + memory_scope mem_scope) { return std::find(capabilities.begin(), capabilities.end(), mem_scope) != capabilities.end(); } From 6d419b9c0b699eb1e4b75f5243e40a410d2fffb6 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Fri, 17 Mar 2023 10:09:01 -0700 Subject: [PATCH 3/4] Addresses code review concerns. Signed-off-by: Maronas, Marcos --- .../.atomic_fence_capabilities.cpp.swp | Bin 0 -> 20480 bytes SYCL/AtomicRef/atomic_fence_capabilities.cpp | 11 +++++++++-- 2 files changed, 9 insertions(+), 2 deletions(-) create mode 100644 SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp diff --git a/SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp b/SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp new file mode 100644 index 0000000000000000000000000000000000000000..76e343abd024913e1ad2f6d28e7c2f16eecb608e GIT binary patch literal 20480 zcmeHPTZkn`8Lmy@C6_fJ*{3yPWx@_U$(%EnB<^zN0+VsFSzULQ*%@V7Hcg-II&(VH zr@N=Cd(KYc;)4r5dqYUT1W_SDamj;9Kvo2OiwIHC2ucD9?n6Y3=1JrCSJgSE&t)bn zj{0CVeAB0@|N5)`ukKYpQ=k6isZa0|)ftB8F2;WJ`6I#g)uZgqm0MY>X+LPu`gWTOTBs+g?T7s)lbv?^0wbB1D1iq?4jA2Y2xAjvHSQv zcW+oqECZGS%YbFTGGH073|Iy%1D1jRBL?#HX7(A_dy}#Ggn1ttc;8~eoeloO1K)2p z;akmH41701`?3sJ1}p=X0n318z%pPNunbrRECZGS%YbFzf53njF!uI)8T-rokU8%E zJNy4@cQW=Y&<4%{Yrwt02Z6Ufz}V}+E5P@F?*bBVfxCg99A@k(paL8L-n@gcH-Kxv zvw#FX4x9uY13m=2eLG`62VMn!2wVl82d)601&#xMyN$750xtv40bd240b0O3a651v z2N2HzSAa)=4+B4YKV#nqS^x*g*5i2p9q(TN-vC13G;kca19%mjd>gm|YysziQ@~N+ z_4hFL1K$=Drb(8Iz?FWI3a1xtiB^*B>S#I1%fVWW z+%8;rS+pi5<{0NwQ~XmDvzaADI>}{G*FU)z2fQn5bnqk#M8;c%%z0fv*p#3Oj@7eW z-4wYSh@H?E&V&ktaU+?hKed{Uuhse-I#l8N+&UH=RWg|ewc3vG^CX)uo3Pl+*Dc3O zzCgtoMUX1?FhMH1N017lbpE!RiAdA)Z)<}l;Yo2o+EBJ4%nnK$%H~}XLAN2S>F^{s z*-Ua%des;VFaN=&peo=7n}U)Z(G+x@>ZY(=h(hq~ITpq_Z+T(tOz@`}yWDFfnUJcj zQ55Hw)o=PmPk#turlPxl2E+6L`drWVFHB$L3%ub)Qcxf3f5HZL9tgbLG2NYU*WP6p6@1O*^G5sE4@qqR|jp70J`c%LS(?37t|% zIZuv`;syU?zPv(9KEZZ^jceVfuwU zM_Scq&yXqGRZQxwFn2DDUv!$7)*}oB<7BF#RGUK8yQmG;5zaVSk2#=IXES=DKo?d7 z>jWh-oN)^5lYn~HV#g9^93~A5V|>OiYO*Xh3};qVP4`o|1(|PX^yEeA=DVo?Wp$Ca zg^2ZfMnWyVx0@pF6O!m_k?1UcXwhSm@kUZ)Ts2~;l4%9&O@V%2&Rk87r53Embji~! z*$D%|J?=M!zs19tCoz)JLW_)~Nl6X=D5|)WVxfbj93p7gpl`&!8oNxG$nAvHI0gO? z)}th3{3Mn+6lE%PStFz#o+Y(C5i>ZJMnnkBKJpQg0!_eT;w-cg+Im1=iKbe@d28JqDP+#RS2vBb0oCfnbc7GWA zQGtFoA0{$K@mq_fI51E9rBc*bc~#T5*sM0k7(?Sfv%0=g<0rA0!X%>x3#l4~nYZ-{ zs?QHyC0pA#Q{%^$mX}@pa?hS#T3&g4_4G$q*7&i!niM&y)-os0ZoI3^@(sxN5yr_& zmw2T?RG3#K_cF|Qa^aZrvPbt#Bf*bJUP(Fyq5c1p*t5R~(EcCWW%XReKHmkN0A_%j zfL}xASHQQ7{0o%EK9&K?fMvikU>UFsSOzQumI2FvWxz6E88}b|=zvMx1NzuN^Z9u` zwvGcnp3<&9sE$#L`Ydp&IY%c_^o%-YP=^!!hvn*$K%G>mQ)^z^>t2zrXL`c=m}7*n z`XDE$6r*kpRVKrCkSe-&2RU?rzW2r9zof8xuS@!9|KG%Y_gUPB)BeBX-~Sr+`_BM$ z2JkW92yh+yeY*4iBCrL_0q@{`{|(?3;A=n#ECWY@L%^SKxBolfD$oYb0S<5r@LSyJ z{}iD6{+ECkfo}p|0lo};3HTiFY2ayK1^6rO_}>Cv04@NBfonLl*alLd23&_Ne*~Tf zlwAkf4y`iFfMvikU>UFsSOx|eIH(1Wy60E7ytpc{UTNBkPc190R~o#I%q^z%N~29q z=NWXd8Vr-1@glKaX_~ZA6j`sd|B_c)2q%@S@beJfDsX!denCH(u9?nX4F4rhIsKpC zI9wJs3;2D-xe&?~LXC_ADYavz@v9`?kMaYxC9le^BLq3xu(BLERcfy@<-u&6eWEZH zUgm}Yyh!2qG2L;W>ZF8ra-Janhu(w8C3tZtUlQZVCb?`TMHKK)SmE0=Nbn`8e3l3eeu literal 0 HcmV?d00001 diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.cpp b/SYCL/AtomicRef/atomic_fence_capabilities.cpp index 16b8cc7bdc..74e8617e40 100644 --- a/SYCL/AtomicRef/atomic_fence_capabilities.cpp +++ b/SYCL/AtomicRef/atomic_fence_capabilities.cpp @@ -12,9 +12,16 @@ using namespace sycl; void checkFenceBehaviour(memory_order order, memory_scope scope) { auto q = queue(); - memory_order order_write = order; + // Both read and write being release or acquire is wrong. In case order is + // release or acquire we need read to be acquire and write to be release. + // If we flip both acquire and release, we will be checking the same case + // (read == acquire, write == release) twice, so we just skip one case and + // flip for the other. + if (order == memory_order::acquire) + return; memory_order order_read = order; - if (memory_order::release == order_read) { + memory_order order_write = order; + if (order == memory_order::release) { order_read = memory_order::acquire; } From 33114c7252030b8e7d73072560ab9c5cc0dd69b7 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Fri, 17 Mar 2023 10:11:54 -0700 Subject: [PATCH 4/4] Removes unnecessary file. Signed-off-by: Maronas, Marcos --- .../.atomic_fence_capabilities.cpp.swp | Bin 20480 -> 0 bytes 1 file changed, 0 insertions(+), 0 deletions(-) delete mode 100644 SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp diff --git a/SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp b/SYCL/AtomicRef/.atomic_fence_capabilities.cpp.swp deleted file mode 100644 index 76e343abd024913e1ad2f6d28e7c2f16eecb608e..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 20480 zcmeHPTZkn`8Lmy@C6_fJ*{3yPWx@_U$(%EnB<^zN0+VsFSzULQ*%@V7Hcg-II&(VH zr@N=Cd(KYc;)4r5dqYUT1W_SDamj;9Kvo2OiwIHC2ucD9?n6Y3=1JrCSJgSE&t)bn zj{0CVeAB0@|N5)`ukKYpQ=k6isZa0|)ftB8F2;WJ`6I#g)uZgqm0MY>X+LPu`gWTOTBs+g?T7s)lbv?^0wbB1D1iq?4jA2Y2xAjvHSQv zcW+oqECZGS%YbFTGGH073|Iy%1D1jRBL?#HX7(A_dy}#Ggn1ttc;8~eoeloO1K)2p z;akmH41701`?3sJ1}p=X0n318z%pPNunbrRECZGS%YbFzf53njF!uI)8T-rokU8%E zJNy4@cQW=Y&<4%{Yrwt02Z6Ufz}V}+E5P@F?*bBVfxCg99A@k(paL8L-n@gcH-Kxv zvw#FX4x9uY13m=2eLG`62VMn!2wVl82d)601&#xMyN$750xtv40bd240b0O3a651v z2N2HzSAa)=4+B4YKV#nqS^x*g*5i2p9q(TN-vC13G;kca19%mjd>gm|YysziQ@~N+ z_4hFL1K$=Drb(8Iz?FWI3a1xtiB^*B>S#I1%fVWW z+%8;rS+pi5<{0NwQ~XmDvzaADI>}{G*FU)z2fQn5bnqk#M8;c%%z0fv*p#3Oj@7eW z-4wYSh@H?E&V&ktaU+?hKed{Uuhse-I#l8N+&UH=RWg|ewc3vG^CX)uo3Pl+*Dc3O zzCgtoMUX1?FhMH1N017lbpE!RiAdA)Z)<}l;Yo2o+EBJ4%nnK$%H~}XLAN2S>F^{s z*-Ua%des;VFaN=&peo=7n}U)Z(G+x@>ZY(=h(hq~ITpq_Z+T(tOz@`}yWDFfnUJcj zQ55Hw)o=PmPk#turlPxl2E+6L`drWVFHB$L3%ub)Qcxf3f5HZL9tgbLG2NYU*WP6p6@1O*^G5sE4@qqR|jp70J`c%LS(?37t|% zIZuv`;syU?zPv(9KEZZ^jceVfuwU zM_Scq&yXqGRZQxwFn2DDUv!$7)*}oB<7BF#RGUK8yQmG;5zaVSk2#=IXES=DKo?d7 z>jWh-oN)^5lYn~HV#g9^93~A5V|>OiYO*Xh3};qVP4`o|1(|PX^yEeA=DVo?Wp$Ca zg^2ZfMnWyVx0@pF6O!m_k?1UcXwhSm@kUZ)Ts2~;l4%9&O@V%2&Rk87r53Embji~! z*$D%|J?=M!zs19tCoz)JLW_)~Nl6X=D5|)WVxfbj93p7gpl`&!8oNxG$nAvHI0gO? z)}th3{3Mn+6lE%PStFz#o+Y(C5i>ZJMnnkBKJpQg0!_eT;w-cg+Im1=iKbe@d28JqDP+#RS2vBb0oCfnbc7GWA zQGtFoA0{$K@mq_fI51E9rBc*bc~#T5*sM0k7(?Sfv%0=g<0rA0!X%>x3#l4~nYZ-{ zs?QHyC0pA#Q{%^$mX}@pa?hS#T3&g4_4G$q*7&i!niM&y)-os0ZoI3^@(sxN5yr_& zmw2T?RG3#K_cF|Qa^aZrvPbt#Bf*bJUP(Fyq5c1p*t5R~(EcCWW%XReKHmkN0A_%j zfL}xASHQQ7{0o%EK9&K?fMvikU>UFsSOzQumI2FvWxz6E88}b|=zvMx1NzuN^Z9u` zwvGcnp3<&9sE$#L`Ydp&IY%c_^o%-YP=^!!hvn*$K%G>mQ)^z^>t2zrXL`c=m}7*n z`XDE$6r*kpRVKrCkSe-&2RU?rzW2r9zof8xuS@!9|KG%Y_gUPB)BeBX-~Sr+`_BM$ z2JkW92yh+yeY*4iBCrL_0q@{`{|(?3;A=n#ECWY@L%^SKxBolfD$oYb0S<5r@LSyJ z{}iD6{+ECkfo}p|0lo};3HTiFY2ayK1^6rO_}>Cv04@NBfonLl*alLd23&_Ne*~Tf zlwAkf4y`iFfMvikU>UFsSOx|eIH(1Wy60E7ytpc{UTNBkPc190R~o#I%q^z%N~29q z=NWXd8Vr-1@glKaX_~ZA6j`sd|B_c)2q%@S@beJfDsX!denCH(u9?nX4F4rhIsKpC zI9wJs3;2D-xe&?~LXC_ADYavz@v9`?kMaYxC9le^BLq3xu(BLERcfy@<-u&6eWEZH zUgm}Yyh!2qG2L;W>ZF8ra-Janhu(w8C3tZtUlQZVCb?`TMHKK)SmE0=Nbn`8e3l3eeu