cuda-oxide 0.4.0

cuda-oxide provides a high-level, rusty wrapper over CUDA. It provides the best safety one can get when working with hardware.
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79

P�U�8(@�4ELF3�q@�44@8@.shstrtab.strtab.symtab.symtab_shndx.nv.info.text.matrixMul_bs32_64bit.nv.info.matrixMul_bs32_64bit.nv.shared.matrixMul_bs32_64bit.nv.constant0.matrixMul_bs32_64bit.text.matrixMul_bs16_64bit.nv.info.matrixMul_bs16_64bit.nv.shared.matrixMul_bs16_64bit.nv.constant0.matrixMul_bs16_64bit.nv.rel.action.shstrtab.strtab.symtab.symtab_shndx.nv.infomatrixMul_bs32_64bit.text.matrixMul_bs32_64bit.nv.info.matrixMul_bs32_64bit.nv.shared.matrixMul_bs32_64bit$___ZZ9matrixMulILi32EmEvPdS0_S0_T0_S1_E2As__170$___ZZ9matrixMulILi32EmEvPdS0_S0_T0_S1_E2Bs__172.nv.constant0.matrixMul_bs32_64bit_parammatrixMul_bs16_64bit.text.matrixMul_bs16_64bit.nv.info.matrixMul_bs16_64bit.nv.shared.matrixMul_bs16_64bit$___ZZ9matrixMulILi16EmEvPdS0_S0_T0_S1_E2As__32$___ZZ9matrixMulILi16EmEvPdS0_S0_T0_S1_E2Bs__34.nv.constant0.matrixMul_bs16_64bit.nv.rel.actionG
�Az
�	2
�,@/	(#			/(#7q0*
@(( �!�!�!�!�!�H�7q0*
@(( �!�!�!�!�!�HK
/" (08 (08 (08,��D���L�����g��� ��"���\'��#���\Q���?W����g�N�"��w�Nw�Og�N�"�B�g�0N��N	g(N� ��
��0NW�0[g�N� ��g�OWH8
��O� ��	G��\	g�L���L� ��7�0[���L�(8� ��\��NwL�B��WH8W��6W@�6� ���(N	w�\WH8�&��7�[7�[� �\� ��W\���\
��N�"��	��O��NG�[� ��g�[��l[7�\�"���
�0[	g�[�0[�`��
G�[g� [g�([� ��	�G\��\�\�"��
��0[�0[��\�@��		�\��\�\����	�\���K� �@�[@�w�Nw�O�"�@�	g�N
��Ng�N� �@�g�Ng�0Ng0[�"�@�g�O
��N��O� �@�	g�(N�(N��N�"�@�g0[��0N
'�0[� �@�	� �\��N	��O� ���\ �\"���\� ��g\7H8�	0[� ���H8
\#���\� �@��g\'�\w�\�����\�\	g\�B��&'���K'7�'�\� ��&& ��	\$G���K� ��%W�$$ �������W\��\4��� � �&]�$]����� �D��M�N�1 �DM� M�N�}@�? 0M��p[@M�R��ȅ�p[N��p[1��PM�`M�  �p[]'@D<N��p[ pM��@*�\�M��p[N�" ���p[�M��M�2��HE  �p[N��p[1��� �M��M��p[]'@D<N��p[�M��@*�\�M�  �p[N�" ���p[ �M�M�2��HE�p[N��p[1���M� M�  �p[]'@D<	N��p[ 0M��@*�\@M��p[
N�" ���p[PM�`M�2��HE  �p[N��p[1��� pM��M��p[]'@D<N��p[�M��@*�\�M�  �p[
N�" ���p[ �M��M�2��HE�p[N��p[1����M��M�  �p[]'@D<N��p[�M����D]������
p[����'�h[G�\�p[��7\"�p[��@�����" ���P�P��������@��P��D���L�����g��� ��	���L'��"���\Q��W��#���\���"�@�g�Nw�Nw�O�"�@�g�Ng�0Ng�N�"��g(N��NG0[�"�@ğg�O
��0NGH8�@��
��O7��\G0[�B����N�(8\� ��g�LGH8�(N� ��G@�6wLW�\� �����L� �\
��N�"����Ng\�'�\�"��G��6	��OGH8�"���0[��l[W�\� ��
�0[	w�[
�[�"����[w�[w� [�&��'�['�[w�([� ��	�7\��\
��0[� ���\�	0[� �\� ��		�\��\�\���@�	�\���K� �@��@@�	w�Nw�O�"�@�
g�N��N
g�O� �@�g�Ng�0N	��0[�"�@�g�N��N��O� �@�g(N��N��0N�"�@��(N�0[
�	0[� �@�� �\��N��O� �@���\w�\7H8� ��	�\"���\�	0[� ��
\wH8#���\�@�@���\'�\w�\�����\�\	w\�B��&'���K'7�'�\� ��&& ��	\$G���K� ��%W�$$ �������g\�G�\7�h[�����\'\4��� � �&]�$�]����� �D���M�N�1 �D�M��M�N�}@�? �M��p[�M�R��ȅ�p[N��p[1���M��M�  �p[]'@D<N��p[ �M��@*�\�M��p[N�" ���p[�M��M�2��HE  �p[N��p[1��� �M��M��p[]'@D<N��p[�M��@*�\�M�  �p[N�"���<�p[�M��
p[^ ��p["�p[������������@�" �������P�P�P��������@��P����P�P�P@9y,��)p�`Mp��
�p��*p���h
Hh2�	�( ��@	( k�@
��@������PH�	�	@4 
&
�.version 7.3
.target sm_52
.address_size 64
3� isible .entry matrixMul_bs16_64bit(
.param .u64"_ ?_0,*1*2*3*�4
)
{
.reg .pred %p<3>;�b32 %r<16�f64 %fd<57%�rd<41>;

	.shaI�.align 8 .b8 _ZZ9��ILi16EmEvPdS0_S0_T0_S1_E2As[2048]G,BG"ld".u�_18, [>0];49414/20424143424�4];
mov.u��5, %ctaid.x6`y;
cvt]0d1,!.R7, %tC,2,7,8,X3,�8;
mul.lo.s�
|T;
shl`440,#e4;
add545)-1'5, �40;
setp.l�X %p1,%�Dwide�"6,4)16�z4�%8,�2d22
2{56, 0d0�;
@%p1 bra $L__BB0_3;
-"32�"r9Vm(25n
�'9,"+2;Q$10��211,k	�?12,��8#3,=Y	g#3,�30%1,53��8��/4,	k#5,>��5�423M�a.to.globalTd�?d19!3�0�93�D6;

%(2:%626,�)40�2�6r'8,�!27]��A6, [*S];
st�a[%r1],$
U'9,�)39�430,$
�731,!/30�7�/31�2��7;
bar.sync D'A%fd8D/4]9bdfma.rn�@�@"11\!+1- 2a.+8c$3,"I�#14e?256e5e.16f$6,#J3F7f/38+#18f.24f$9,#J1#20fO512] 1f-32f{#Jf9F3f?640f4f.40f$5,#J2F6f/76�$27f��#J5F9f/89�#30f-56fa#J8F2fO1024g3g6�434,#K1F5g/15#36g7437,#K4F8�#39g-80gq#K7F4�/404�443,#K0F?1534.96g$6,#K4O1664g.10449,$L5?1795=112h452,$Lh5?1925>120h$6,$L2k�'9,7�
$16�
e�7p2,#5
2
23j�
*d2�'3,"
q��

C
XA4��Z
�$7,E�'8,6!37���%8]kret;

}�/32�
/32�*1*2*3*�%6105&� /32�1[81�G*BG�/32�/32�/32�/32�/32��5�/24�L/32�5�?104�1�n8�
/32��/32���:
�1��j?104�>�G?512�Pq�A�B�
�B=
�B��B��MO2304g6�:?256�Q!28�MO3072�PO3328�QO3584�QO3840�5�G	O4096h7h458,$
5
/43�6
=136h461,$L6/46}6.14464,$L6/486�h$7,$L
&68C%69D470,$L7?537%72h473,$L7O5632h>176h$6,$L7O5888h.84h$9,$L86P!-92h482,$Lh8�8>200h$5,$L8/66
%87h488,$L8/69
92
491,$L97�h2
494,$L9/74
92 
497,$L&98�
%99�
D100,%M$10/79#
5102k)&P0.-/32.1.1.��
�ret;

}