From b6d639a2eb8e4466c34a0f154ce16d45223b6883 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Sun, 9 Feb 2025 12:54:40 -0600 Subject: [PATCH 01/11] Export _md_local_accessor Cython struct, mapped to C struct MDLocalAccessor --- dpctl/_backend.pxd | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 1f27a1f540..9b41499160 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -362,6 +362,12 @@ cdef extern from "syclinterface/dpctl_sycl_kernel_bundle_interface.h": cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": + ctypedef struct _md_local_accessor 'MDLocalAccessor': + size_t ndim + _arg_data_type dpctl_type_id + size_t dim0 + size_t dim1 + size_t dim2 cdef bool DPCTLQueue_AreEq(const DPCTLSyclQueueRef QRef1, const DPCTLSyclQueueRef QRef2) cdef DPCTLSyclQueueRef DPCTLQueue_Create( From fd01ba5654683a3076a1dcbc8a277f11b79aae1e Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Sun, 9 Feb 2025 12:56:17 -0600 Subject: [PATCH 02/11] Define LocalAccessor type to use to specify local accessor kernel arguments LocalAccessor(ndim, elemental_type_str, dim0, dim1, dim2) The elemental type can be one of the following: "i1", "u1", "i2", "u2", "i4", "u4", "i8", "u8", "f4", "f8" --- dpctl/_sycl_queue.pyx | 45 +++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 45 insertions(+) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 94527506ef..e4a32f7fa7 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -59,6 +59,7 @@ from ._backend cimport ( # noqa: E211 DPCTLWorkGroupMemory_Delete, _arg_data_type, _backend_type, + _md_local_accessor, _queue_property_type, ) from .memory._memory cimport _Memory @@ -125,6 +126,47 @@ cdef class kernel_arg_type_attribute: return self.attr_value +cdef class LocalAccessor: + cdef _md_local_accessor lacc + + def __cinit__(self, size_t ndim, str type, size_t dim0, size_t dim1, size_t dim2): + self.lacc.ndim = ndim + self.lacc.dim0 = dim0 + self.lacc.dim1 = dim1 + self.lacc.dim2 = dim2 + + if ndim < 1 or ndim > 3: + raise ValueError + if type == 'i1': + self.lacc.dpctl_type_id = _arg_data_type._INT8_T + elif type == 'u1': + self.lacc.dpctl_type_id = _arg_data_type._UINT8_T + elif type == 'i2': + self.lacc.dpctl_type_id = _arg_data_type._INT16_T + elif type == 'u2': + self.lacc.dpctl_type_id = _arg_data_type._UINT16_T + elif type == 'i4': + self.lacc.dpctl_type_id = _arg_data_type._INT32_T + elif type == 'u4': + self.lacc.dpctl_type_id = _arg_data_type._UINT32_T + elif type == 'i8': + self.lacc.dpctl_type_id = _arg_data_type._INT64_T + elif type == 'u8': + self.lacc.dpctl_type_id = _arg_data_type._UINT64_T + elif type == 'f4': + self.lacc.dpctl_type_id = _arg_data_type._FLOAT + elif type == 'f8': + self.lacc.dpctl_type_id = _arg_data_type._DOUBLE + else: + raise ValueError(f"Unrecornigzed type value: '{type}'") + + def __repr__(self): + return "LocalAccessor(" + self.ndim + ")" + + cdef size_t addressof(self): + return &self.lacc + + cdef class _kernel_arg_type: """ An enumeration of supported kernel argument types in @@ -865,6 +907,9 @@ cdef class SyclQueue(_SyclQueue): elif isinstance(arg, WorkGroupMemory): kargs[idx] = (arg._ref) kargty[idx] = _arg_data_type._WORK_GROUP_MEMORY + elif isinstance(arg, LocalAccessor): + kargs[idx] = ((arg).addressof()) + kargty[idx] = _arg_data_type._LOCAL_ACCESSOR else: ret = -1 return ret From c0019d54f81f5a83ca6ba43437d5be32ee1a63f2 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Sun, 9 Feb 2025 13:19:53 -0600 Subject: [PATCH 03/11] Add Python tests for kernel submission with LocalAccessor --- .../local_accessor_kernel_fp64.spv | Bin 0 -> 5696 bytes .../local_accessor_kernel_inttys_fp32.spv | Bin 0 -> 12908 bytes dpctl/tests/test_sycl_kernel_submit.py | 41 ++++++++++++++++++ 3 files changed, 41 insertions(+) create mode 100644 dpctl/tests/input_files/local_accessor_kernel_fp64.spv create mode 100644 dpctl/tests/input_files/local_accessor_kernel_inttys_fp32.spv diff --git a/dpctl/tests/input_files/local_accessor_kernel_fp64.spv b/dpctl/tests/input_files/local_accessor_kernel_fp64.spv new file mode 100644 index 0000000000000000000000000000000000000000..ffc220268a8f6da9ad169350ad2820ea3e66e859 GIT binary patch literal 5696 zcmai$33n4!6vrntfZaf4b3tuE0WG2+f>nW9gi zp)87m3WEAE{0x2x&r#1o{r_I(Hq+^FoO5pPe&5~RODo3A9PcF)-UM%)_p4X>Zt})> zNs-{O0Uj6NssL99cvgUC2l(2QlHYIBJg-VI+c#|STi0z~(dOs!#cVE{mR*%M-pSXj z+T^GA=F^=;Kh@XQzYATpPD}`|&!@A?+vfEZyFBGqp-emn6TM1B^}HJe{+D^VE@acaetX;6)~*$@u5kGe9Ze`25oPQ5^$!7X?gzPH4o;bTwj(x@c4p+y)IGZ4rGm@S&2S?kL>j(_7 zJ2~2JxCV>Y;?*jq=iTdU|9>CY!xrrlK2sf?bTqp^J;cuyUFY~(x0^BXV|U2nvvU?7 zn)tW7cu5zp9KYJ_Tul5CduMUkPm2pp+}nb<@a4EQZck(4j@ng=T`7Ny4NdGBf)?F8 zcw*n+_!_s{F|mhpw*9u)+y#pbP3&7;tfY%o&O50$DtuPbyLkj(<8fa~XOFlGqwb8% zlzXkmy@|0e@7(Y$!vmihuR%D*XL$d@F6`3+*$*gjdi99tIRf6Zk-7?Rsbu}mu0)r7 zD!gUG@)~z1W6zGLtnu)^r zVTZ0>{`&XE$tEjiz3A!kSS8Fk^np07!s7+V;nxQ|eSqH>@c6*D3zIK=yt3@4hVvCjlNG2qV$bC&7P2R!=6q_;lO7vuSD3iHG_o>-M3 z9)I3$^y~+{yhjjxO^YO>2HrhBdc0VDrN?^m98Irz&;L~QlY|?U6+ZkH@UAaWkC-#R zX`EH6-BVvncY)tfbbMyZ?>0F_ZSRl_9iQ3WDVmz2ws%Q}j_}K`ox*^)M}nBdY0gK^?>^faF0ztS9F8G&b(hV_16i= z2mUWboh(d0!GCB&nD>K%c>;R;fI$DkI~_SZB;Y+j$L|(8J7b|}^qljz@`6XtZ#QQ^ zTb&P!rcUhXFZ-}afZS@$DF*Vz0&+U6e!%k^i7oQ1Y{;2Bksp!`{^Uk%-l=e%P zd0aFz!iO0Az9QotI}*s2iRNr{mQRyt&e&1AV`&4O!QLb)bX zLuu#P-GqQQgg!kf%pK#sSS{dQU65_P@H#=8V6A{2!><&q5a0_wAUSW#1_AzFqQcRf zvr@h0o!KNn$1H(;@@W^|EI?;F|CDIXrqyr3$O*9lwP7eu4uvmW@O z<7~G}j*c4G7tTlth_PK@{vD#@{5vH_hoAX(i6#a*+tajYi5)iJ=$X7l z;WOV7Z##ZWba=<#6@Oe{Il#Xocw0bxu(fG6D(BU0_FVIDI{C&yM*?c|_9k=5j zN{)_?9m5A1^WrVG9sfu)IdD#R$3GU2KXv1e9kQtZCz6?ec*j4L9DnTb#}1kK+m3%G zS$M}AwfoHHOTkG2^M2F|KTYO$%^McO`b_s|N zws{pqoHZoz1IXblkiKBuB?bzwkk3{jz!Ok&GNTBRsFY;_
yVfH+?XGt*PTLjrmVzft&fh(9CxY``;Q_p)eYoC}T5 z_0HfgoHH+0&3X_ozSGy-E?;X!H5A7 zX8@Jp3T~(fiZd+A>WYdBtD=CiDLUgmGtTmx5f{Mkd-YD$t(Vu)`^(AwpXHu=&b|Nl z>aopxJ6guEBmF~% zFPS?wetJf2H7dz+*(RGx)+~D-(meIVrTgr+blvj7#r>m0{e#V=OO9UDH$RJ*!zYS4 zqZ6}gqL{Hx%w-eBjCW#gnQ2m26UD5U&uP~3 z`@QO-9!}d6Ctq`HWMFh{^RP7ogX4>a77h+C>l<7&w03xT-}u1r&?2rOkM-8n*vma{+BB}yM-gb; zU7NsBdxIOfJ|lm2EAm%zDDum> z$S*ekOY(g6JYOgOw48Gl`MY=*dF5y16`S|PMPBtz-sw3{EAno_)yO@S^^sd_?p+Yi zNZs69=KA!U+ZDOnH3z>Vx6eZ47MuG;d9He%t5bKKcX(^B*4ldZf-Ze}=JV3pyX#p9 zKDM73!l->l=5tdK-{D->3fD}}=2Wj~^DkW@pY|nwErplyA=uXxc5y3adj8y1)HKj6 zJ$o$4&;JN!&nV7hewSX$PDYP?#8JyB3I2OxmL++_k!N|^Ue5->*5dD*+*sYE*Rl)H zV_k|b!5%e=TQT))W7qn6c6pax%dSq=i6i$l3I2OxZbS9;>i7U+g{I}1)IC| z+}5+_y7XGM9fXQ?Dc%Wt^eK**o!j<$wktTzy<3-F%MMJ|rMS19qn3#y_hD^&JzES; z&tyrLUd#I0>+0D`@SgPL_m$6I$5+c%;fp@RvDaYRUeDHo(>~5d=TFAl<++8tsr>D$ zd{$e((eYE1^W}*Smh>k|`jd%1ljqTQ&Y>O-Lw!24+;*tT zgN6!2eI>KJaLnPC-t`TYEZg8;bBY|hgY~JK!!t?sm!V%)>-guQJ0j;kEw#%w4~_B* zTPyFuh0T4eZ9UT3r&Zh5EKqku?lEj^osWHsdgZ7i>OZH@!!Pd~q06IGmj{j=$Y*Z% zbZ($~cYYIv?)j=4vn_v5V>WI^VfPNyw`O*&`Ycv#!)&Z^`eu=Fhr+fUd1irKcLuZh zJrDH(X7i|<=S*hv7^lyeo0(nXU8&!~>>72~+{)}4Xej(UBTDE$2B)aR>UB8Ig^~UMj5$`y&`MopuT*EvCUAy|Zg&t?NzR+Ex|2%Z- zY1KG6yA0huv}=q0t|{4PvG&Ho7P;<5cW>94`@TXqhx!BP?iKzA3tj&~`0dsHmC*f` zO?C5!Zmy|`?s~uV^vhGv?KFgO=>AykE1pfcascNqSNL4Yv7GVS?`b;UmU$jIT;q46 z-++hkB`rt&{&k;L-&uA%xy62O>vKe`-}u&3j@=2HK1amv47T5L>`SrfbKFR_S>{Ti^j^|{WvUW;hAcKLWHA}7t|9jO}=b;lhXhOHNQhMPMO369HE z@6val=Q4ug^W?k1Zx5fL>&QEcc^$g%cHifosro2#7UEjx_1JxnyajQ6)>F%|`z+vW zeW$#&@OgL2I@b7{zpd~E$2Xw6zr8e&cy`XmuFql5=F?`J&w1=~0XA!J5APz@za5+F z`IE0VPFs2Xh1g=fzkRX(9bnfR@BYSVi}ltoz6hJ|oYd|+;hos4-UL5QG(-sgQgt61KsbBs6qW!g+j{_EIceai7~V0XRo?r)s7SRWk!CbpF0 z^WeVs^BwGdzuT`dawN{k`|?aZPv`P|^|Bsw`@TMa68CilIK8hAqhE={eW`y4`5iPQV~XyGg0*T=Byi~IUGxaz*H!miKZdAUYgoLAh}C$L#V`My4h&Gr1r z*Bhs;y#7y8`|?aZPv`P|9iDfc z<=@otopb~?vELE%5!WtYj^CrBz@C4c>tb-#duj=GeU6Ad8eH|AbPRTVj)*-LT=kuF z9Cm$<_+~gBY>o!vcZquVP5`^6f#_3@?}w#ed(`jUt5^T`YasFMu@q4-8_mSi3`p-h#LmYj+18kqhd%i1~FG93OufffC z7QW!-26SuG9vrt8a|OqvFZYf-7lXZ%$a4u;TgvfE(ao*Tb=LJRM7y=i$9E%g(p=t= zx-n6A+|hfm^&(}CUxv+*a(pAYe3t{4Bc9iLnIqQ+!70Z-gkHt*E3oTx1aGedS8@Eq z*!4Lg_9Ngbj(-%pK1ba5$H3-jAaY7Qd>;q9rh({FkGsDLY>)aqKlLh(e**jFIQ~ie zDaSvBZvPF$J=Nv-`UvSf!tWkS#+*-^P9FM--JMvr)_D&+t4Pb34$8SV8 zx4z)`O^H$x9KQ>@K1ba5x4`CTAaY7QeBTDU zrh({F_x$YdJBU5%_x#lNLaeWjM9<#^&qgMCkKc_i<@i15cOnhMJ=Nv-mzX_I=ivB# zU^(s?--~=1364JimgD;FNAyKs-vjHv6OrTU`oDs>hdBEDKG;5u_k8zceh|?fy#^Pp z!J703r`_AN!RhG3y@JyZfxUys^DtQ3!ASf&oJYXs)@N<vP0?{t9f41|nzF!}n{jYZ{0?^|;&LfbCJg=ciu9;ooAPjZDPh-{DI+ z{O{=Y-$2|`T@F9V?0Gr|hyN2i7x9e$1NjLO9R4q`9M<=HL|^pv-(dajBZt-X{}gc# zarF5=VEZ&aIQ+k0?VfKP3GP{gHR%ryySHnD!_kL(1&9A1>>Wg&KY+ER9DW+z-1>sU ze+0`p*Lnx)#zkGhyFX#;Mamq02AktAh`imB`Pl;hnfW=y^T>KS9Ik^?4o^a_;_zhb z`W(T_DcaDjuYm+Fx5Tc`5wTOzC!?2Rx5BQ^5%)O_-5d>Uaz;IT)4{F@pStH~e>1@L zsNeHb_n+t8)3b>4+Xma_IQ%^PDTkkrZvPE(SdY3K?qT*korA;Mg5_|=`~qZ)!e}E>Qx+n752?>d_Vjt$6u|E#GdMMd^WS^ z=^PwC04&En=%|-L-on?jepo4+7h#@xk$f!P z!OcUk=?iXJlQn7&j$4blg5%Mbdq=;`Ier+rx%CCd4+m=xjvoP*ljib{ z)QyR{gO~HM^>S8ajvs~Hv4~k-FJxX^;3JurAfDIJ%#rJOaLVx$(5pDU6uUl0aQsB{ zDvrMnyS_B`B=jndpNw6fBkue4=;ml(ljG{)I|b~T@Tq%#_ID~`kNQ17_0tf~uZ~2| zZ@{)Wj=vFq%JDa$+kfn-F2|2y_B@?~<8KDbanHDk99#ITVHsGC>+36g(bsZx{q7^j z)m?iW;vVAY^K`I%8t+{!VD3k>N3X-=(7pn(rt>)u^_7UR+J}hQlX|REdvM;`%pIJM zKCLA%sRX`hDR03w<5_2f*J$T%+z@ z-$z`d{vh)Y5ZBm?vDRW8{)TBY-kMUIYpgALF86qC?(4~}VU8X@Q()J87HnTXWRAWb z0;hdFT Date: Sun, 16 Feb 2025 18:04:13 -0800 Subject: [PATCH 04/11] Fix typo in `_md_local_accessor` error message --- dpctl/_sycl_queue.pyx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index e4a32f7fa7..e6ffb1d5d4 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -158,7 +158,7 @@ cdef class LocalAccessor: elif type == 'f8': self.lacc.dpctl_type_id = _arg_data_type._DOUBLE else: - raise ValueError(f"Unrecornigzed type value: '{type}'") + raise ValueError(f"Unrecognized type value: '{type}'") def __repr__(self): return "LocalAccessor(" + self.ndim + ")" From 1f94e7f244754b9caa262acdc46ffcd4a0f784b9 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Wed, 26 Feb 2025 00:19:01 -0800 Subject: [PATCH 05/11] type->dtype in LocalAccessor initialization Fixes conflict with built-in `type` --- dpctl/_sycl_queue.pyx | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index e6ffb1d5d4..ad8d509989 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -129,7 +129,7 @@ cdef class kernel_arg_type_attribute: cdef class LocalAccessor: cdef _md_local_accessor lacc - def __cinit__(self, size_t ndim, str type, size_t dim0, size_t dim1, size_t dim2): + def __cinit__(self, size_t ndim, str dtype, size_t dim0, size_t dim1, size_t dim2): self.lacc.ndim = ndim self.lacc.dim0 = dim0 self.lacc.dim1 = dim1 @@ -137,28 +137,28 @@ cdef class LocalAccessor: if ndim < 1 or ndim > 3: raise ValueError - if type == 'i1': + if dtype == 'i1': self.lacc.dpctl_type_id = _arg_data_type._INT8_T - elif type == 'u1': + elif dtype == 'u1': self.lacc.dpctl_type_id = _arg_data_type._UINT8_T - elif type == 'i2': + elif dtype == 'i2': self.lacc.dpctl_type_id = _arg_data_type._INT16_T - elif type == 'u2': + elif dtype == 'u2': self.lacc.dpctl_type_id = _arg_data_type._UINT16_T - elif type == 'i4': + elif dtype == 'i4': self.lacc.dpctl_type_id = _arg_data_type._INT32_T - elif type == 'u4': + elif dtype == 'u4': self.lacc.dpctl_type_id = _arg_data_type._UINT32_T - elif type == 'i8': + elif dtype == 'i8': self.lacc.dpctl_type_id = _arg_data_type._INT64_T - elif type == 'u8': + elif dtype == 'u8': self.lacc.dpctl_type_id = _arg_data_type._UINT64_T - elif type == 'f4': + elif dtype == 'f4': self.lacc.dpctl_type_id = _arg_data_type._FLOAT - elif type == 'f8': + elif dtype == 'f8': self.lacc.dpctl_type_id = _arg_data_type._DOUBLE else: - raise ValueError(f"Unrecognized type value: '{type}'") + raise ValueError(f"Unrecognized type value: '{dtype}'") def __repr__(self): return "LocalAccessor(" + self.ndim + ")" From 76fdd08fff1cfb4dbc49f0c24bdb582392e7243d Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 27 Feb 2025 10:00:50 -0800 Subject: [PATCH 06/11] Add docstrings to LocalAccessor Also improve messages in errors --- dpctl/_sycl_queue.pyx | 41 ++++++++++++++++++++++++++++++++++++++++- 1 file changed, 40 insertions(+), 1 deletion(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index ad8d509989..472f105aea 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -127,6 +127,41 @@ cdef class kernel_arg_type_attribute: cdef class LocalAccessor: + """ + LocalAccessor(ndim, dtype, dim0, dim1, dim2) + + Python class for specifying the dimensionality and type of a + ``sycl::local_accessor``, to be used as a kernel argument type. + + Args: + ndim (size_t): + number of dimensions. + Can be between one and three. + dtype (str): + the data type of the local memory. + The permitted values are + + `'i1'`, `'i2'`, `'i4'`, `'i8'`: + signed integral types int8_t, int16_t, int32_t, int64_t + `'u1'`, `'u2'`, `'u4'`, `'u8'` + unsigned integral types uint8_t, uint16_t, uint32_t, + uint64_t + `'f4'`, `'f8'`, + single- and double-precision floating-point types float and + double + dim0 (size_t): + Size of the first dimension. + dim1 (size_t): + Size of the second dimension. + dim2 (size_t): + Size of the third dimension. + + Raises: + ValueError: + If the given dimension is not between one and three. + ValueError: + If the dtype string is unrecognized. + """ cdef _md_local_accessor lacc def __cinit__(self, size_t ndim, str dtype, size_t dim0, size_t dim1, size_t dim2): @@ -136,7 +171,7 @@ cdef class LocalAccessor: self.lacc.dim2 = dim2 if ndim < 1 or ndim > 3: - raise ValueError + raise ValueError("LocalAccessor must have dimension between one and three") if dtype == 'i1': self.lacc.dpctl_type_id = _arg_data_type._INT8_T elif dtype == 'u1': @@ -164,6 +199,10 @@ cdef class LocalAccessor: return "LocalAccessor(" + self.ndim + ")" cdef size_t addressof(self): + """ + Returns the address of the _md_local_accessor for this LocalAccessor + cast to ``size_t``. + """ return &self.lacc From 5c07e441b55d1f8a158e9282fe9ed833d95c2557 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Thu, 27 Feb 2025 10:01:34 -0800 Subject: [PATCH 07/11] Add a comment documenting the origin of spirv files in test_sycl_kernel_submit.py --- dpctl/tests/test_sycl_kernel_submit.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index bffca34e5b..e96567d4fe 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -288,6 +288,9 @@ def get_spirv_abspath(fn): return spirv_file +# the process for generating the .spv files in this test is documented in +# libsyclinterface/tests/test_sycl_queue_submit_local_accessor_arg.cpp +# in a comment starting on line 123 def test_submit_local_accessor_arg(): try: q = dpctl.SyclQueue("level_zero") From 6cdb80c74248ef65737442b4ea981fd28c9a5603 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 7 Mar 2025 13:21:59 -0800 Subject: [PATCH 08/11] Add LocalAccessor to dpctl namespace --- dpctl/__init__.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/dpctl/__init__.py b/dpctl/__init__.py index 1d9b7209e4..e4dd710ade 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -48,6 +48,7 @@ from ._sycl_event import SyclEvent from ._sycl_platform import SyclPlatform, get_platforms, lsplatform from ._sycl_queue import ( + LocalAccessor, SyclKernelInvalidRangeError, SyclKernelSubmitError, SyclQueue, @@ -102,6 +103,7 @@ "SyclKernelSubmitError", "SyclQueueCreationError", "WorkGroupMemory", + "LocalAccessor", ] __all__ += [ "get_device_cached_queue", From 5fc86c63c50ae9bfc836e61678907cecb15cdec5 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Fri, 7 Mar 2025 13:22:10 -0800 Subject: [PATCH 09/11] Fix __repr__ in LocalAccessor --- dpctl/_sycl_queue.pyx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 472f105aea..13cecd18b9 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -196,7 +196,7 @@ cdef class LocalAccessor: raise ValueError(f"Unrecognized type value: '{dtype}'") def __repr__(self): - return "LocalAccessor(" + self.ndim + ")" + return f"LocalAccessor({self.lacc.ndim})" cdef size_t addressof(self): """ From b5887dfd7a3a19e823db201a2cffc755114b71c0 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sat, 8 Mar 2025 00:24:05 -0800 Subject: [PATCH 10/11] LocalAccessor takes a sequence of non-negative integers instead of separate integer arguments Update test for LocalAccessor kernel submission --- dpctl/_sycl_queue.pyx | 43 ++++++++++++++++---------- dpctl/tests/test_sycl_kernel_submit.py | 2 +- 2 files changed, 27 insertions(+), 18 deletions(-) diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 13cecd18b9..ad44e8faa2 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -128,15 +128,12 @@ cdef class kernel_arg_type_attribute: cdef class LocalAccessor: """ - LocalAccessor(ndim, dtype, dim0, dim1, dim2) + LocalAccessor(dtype, shape) Python class for specifying the dimensionality and type of a ``sycl::local_accessor``, to be used as a kernel argument type. Args: - ndim (size_t): - number of dimensions. - Can be between one and three. dtype (str): the data type of the local memory. The permitted values are @@ -149,29 +146,41 @@ cdef class LocalAccessor: `'f4'`, `'f8'`, single- and double-precision floating-point types float and double - dim0 (size_t): - Size of the first dimension. - dim1 (size_t): - Size of the second dimension. - dim2 (size_t): - Size of the third dimension. + shape (tuple, list): + Size of LocalAccessor dimensions. Dimension of the LocalAccessor is + determined by the length of the tuple. Must be of length 1, 2, or 3, + and contain only non-negative integers. Raises: + TypeError: + If the given shape is not a tuple or list. ValueError: - If the given dimension is not between one and three. + If the given shape sequence is not between one and three elements long. + TypeError: + If the shape is not a sequence of integers. + ValueError: + If the shape contains a negative integer. ValueError: If the dtype string is unrecognized. """ cdef _md_local_accessor lacc - def __cinit__(self, size_t ndim, str dtype, size_t dim0, size_t dim1, size_t dim2): + def __cinit__(self, str dtype, shape): + if not isinstance(shape, (list, tuple)): + raise TypeError(f"`shape` must be a list or tuple, got {type(shape)}") + ndim = len(shape) + if ndim < 1 or ndim > 3: + raise ValueError("LocalAccessor must have dimension between one and three") + for s in shape: + if not isinstance(s, numbers.Integral): + raise TypeError("LocalAccessor shape must be a sequence of integers") + if s < 0: + raise ValueError("LocalAccessor dimensions must be non-negative") self.lacc.ndim = ndim - self.lacc.dim0 = dim0 - self.lacc.dim1 = dim1 - self.lacc.dim2 = dim2 + self.lacc.dim0 = shape[0] + self.lacc.dim1 = shape[1] if ndim > 1 else 1 + self.lacc.dim2 = shape[2] if ndim > 2 else 1 - if ndim < 1 or ndim > 3: - raise ValueError("LocalAccessor must have dimension between one and three") if dtype == 'i1': self.lacc.dpctl_type_id = _arg_data_type._INT8_T elif dtype == 'u1': diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index e96567d4fe..a977a4f7df 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -308,7 +308,7 @@ def test_submit_local_accessor_arg(): try: e = q.submit( krn, - [x.usm_data, dpctl._sycl_queue.LocalAccessor(1, "i8", lws, 1, 1)], + [x.usm_data, dpctl.LocalAccessor("i8", (lws, 1, 1))], [ gws, ], From 1d3453ebf11dcde30e0e2fba5c9106938baccb6e Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Mon, 10 Mar 2025 17:32:27 -0700 Subject: [PATCH 11/11] Address review comments on test_submit_local_accessor_arg Also slips in a fix to the shape of the local_accessor --- dpctl/tests/test_sycl_kernel_submit.py | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index a977a4f7df..e46c4f1760 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -308,13 +308,9 @@ def test_submit_local_accessor_arg(): try: e = q.submit( krn, - [x.usm_data, dpctl.LocalAccessor("i8", (lws, 1, 1))], - [ - gws, - ], - [ - lws, - ], + [x.usm_data, dpctl.LocalAccessor("i8", (lws,))], + [gws], + [lws], ) e.wait() except dpctl._sycl_queue.SyclKernelSubmitError: