1
+ // blake32 kernel.
2
+ // Author: CryptoGraphics ( CrGraphics@protonmail.com )
3
+
4
+ #define rotr32 (a , w , c ) \
5
+ { \
6
+ a = ( w >> c ) | ( w << ( 32 - c ) ); \
7
+ }
8
+
9
+ #define blake32GS (a , b , c , d , x , y , mx , my ) \
10
+ { \
11
+ v[a] += (mx ^ c_u256[y]) + v[b]; \
12
+ v[d] ^= v[a]; \
13
+ rotr32(v[d], v[d], 16U); \
14
+ v[c] += v[d]; \
15
+ v[b] ^= v[c]; \
16
+ rotr32(v[b], v[b], 12U); \
17
+ \
18
+ v[a] += (my ^ c_u256[x]) + v[b]; \
19
+ v[d] ^= v[a]; \
20
+ rotr32(v[d], v[d], 8U); \
21
+ v[c] += v[d]; \
22
+ v[b] ^= v[c]; \
23
+ rotr32(v[b], v[b], 7U); \
24
+ }
25
+
26
+ #define byteSwapU32 (ret , val ) \
27
+ { \
28
+ val = ((val << 8U) & 0xFF00FF00U ) | ((val >> 8U) & 0xFF00FFU ); \
29
+ ret = (val << 16U) | (val >> 16U); \
30
+ }
31
+
32
+ typedef union {
33
+ uint4 h4 [2 ];
34
+ ulong4 h8 ;
35
+ } hash_t ;
36
+
37
+ __attribute__((reqd_work_group_size (256 , 1 , 1 )))
38
+ __kernel void blake32 (__global uint * hashes ,
39
+ const uint uH0 , const uint uH1 , const uint uH2 , const uint uH3 ,
40
+ const uint uH4 , const uint uH5 , const uint uH6 , const uint uH7 ,
41
+ const uint in16 , const uint in17 , const uint in18 , const uint firstNonce )
42
+ {
43
+ int gid = get_global_id (0 );
44
+
45
+ __global hash_t * hash = (__global hash_t * )(hashes + (8 * (get_global_id (0 ))));
46
+ uint nonce = firstNonce + (uint )gid ;
47
+
48
+
49
+ const uint c_u256 [16 ] = {
50
+ 0x243F6A88U , 0x85A308D3U ,
51
+ 0x13198A2EU , 0x03707344U ,
52
+ 0xA4093822U , 0x299F31D0U ,
53
+ 0x082EFA98U , 0xEC4E6C89U ,
54
+ 0x452821E6U , 0x38D01377U ,
55
+ 0xBE5466CFU , 0x34E90C6CU ,
56
+ 0xC0AC29B7U , 0xC97C50DDU ,
57
+ 0x3F84D5B5U , 0xB5470917U
58
+ };
59
+
60
+ uint h [8 ];
61
+ uint v [16 ];
62
+
63
+ h [0 ]= uH0 ;
64
+ h [1 ]= uH1 ;
65
+ h [2 ]= uH2 ;
66
+ h [3 ]= uH3 ;
67
+ h [4 ]= uH4 ;
68
+ h [5 ]= uH5 ;
69
+ h [6 ]= uH6 ;
70
+ h [7 ]= uH7 ;
71
+
72
+ for (int i = 0 ; i < 8 ; ++ i )
73
+ v [i ] = h [i ];
74
+
75
+ v [8 ] = 0x243F6A88U ;
76
+ v [9 ] = 0x85A308D3U ;
77
+ v [10 ] = 0x13198A2EU ;
78
+ v [11 ] = 0x03707344U ;
79
+ v [12 ] = 0xA4093822U ^ 640U ;
80
+ v [13 ] = 0x299F31D0U ^ 640U ;
81
+ v [14 ] = 0x082EFA98U ;
82
+ v [15 ] = 0xEC4E6C89U ;
83
+
84
+ // { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
85
+ blake32GS (0 , 4 , 0x8 , 0xC , 0 , 1 , in16 , in17 );
86
+ blake32GS (1 , 5 , 0x9 , 0xD , 2 , 3 , in18 , nonce );
87
+ blake32GS (2 , 6 , 0xA , 0xE , 4 , 5 , 0x80000000U , 0U );
88
+ blake32GS (3 , 7 , 0xB , 0xF , 6 , 7 , 0U , 0U );
89
+ blake32GS (0 , 5 , 0xA , 0xF , 8 , 9 , 0U , 0U );
90
+ blake32GS (1 , 6 , 0xB , 0xC , 10 , 11 , 0U , 0U );
91
+ blake32GS (2 , 7 , 0x8 , 0xD , 12 , 13 , 0U , 1U );
92
+ blake32GS (3 , 4 , 0x9 , 0xE , 14 , 15 , 0U , 640U );
93
+
94
+ // { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
95
+ blake32GS (0 , 4 , 0x8 , 0xC , 14 , 10 , 0U , 0U );
96
+ blake32GS (1 , 5 , 0x9 , 0xD , 4 , 8 , 0x80000000 , 0U );
97
+ blake32GS (2 , 6 , 0xA , 0xE , 9 , 15 , 0U , 640U );
98
+ blake32GS (3 , 7 , 0xB , 0xF , 13 , 6 , 1U , 0U );
99
+ blake32GS (0 , 5 , 0xA , 0xF , 1 , 12 , in17 , 0U );
100
+ blake32GS (1 , 6 , 0xB , 0xC , 0 , 2 , in16 , in18 );
101
+ blake32GS (2 , 7 , 0x8 , 0xD , 11 , 7 , 0U , 0U );
102
+ blake32GS (3 , 4 , 0x9 , 0xE , 5 , 3 , 0U , nonce );
103
+
104
+ // { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
105
+ blake32GS (0 , 4 , 0x8 , 0xC , 11 , 8 , 0U , 0U );
106
+ blake32GS (1 , 5 , 0x9 , 0xD , 12 , 0 , 0U , in16 );
107
+ blake32GS (2 , 6 , 0xA , 0xE , 5 , 2 , 0U , in18 );
108
+ blake32GS (3 , 7 , 0xB , 0xF , 15 , 13 , 640U , 1U );
109
+ blake32GS (0 , 5 , 0xA , 0xF , 10 , 14 , 0U , 0U );
110
+ blake32GS (1 , 6 , 0xB , 0xC , 3 , 6 , nonce , 0U );
111
+ blake32GS (2 , 7 , 0x8 , 0xD , 7 , 1 , 0U , in17 );
112
+ blake32GS (3 , 4 , 0x9 , 0xE , 9 , 4 , 0U , 0x80000000U );
113
+
114
+ // { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
115
+ blake32GS (0 , 4 , 0x8 , 0xC , 7 , 9 , 0U , 0U );
116
+ blake32GS (1 , 5 , 0x9 , 0xD , 3 , 1 , nonce , in17 );
117
+ blake32GS (2 , 6 , 0xA , 0xE , 13 , 12 , 1U , 0U );
118
+ blake32GS (3 , 7 , 0xB , 0xF , 11 , 14 , 0U , 0U );
119
+ blake32GS (0 , 5 , 0xA , 0xF , 2 , 6 , in18 , 0U );
120
+ blake32GS (1 , 6 , 0xB , 0xC , 5 , 10 , 0U , 0U );
121
+ blake32GS (2 , 7 , 0x8 , 0xD , 4 , 0 , 0x80000000U , in16 );
122
+ blake32GS (3 , 4 , 0x9 , 0xE , 15 , 8 , 640U , 0U );
123
+
124
+ // { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
125
+ blake32GS (0 , 4 , 0x8 , 0xC , 9 , 0 , 0U , in16 );
126
+ blake32GS (1 , 5 , 0x9 , 0xD , 5 , 7 , 0U , 0U );
127
+ blake32GS (2 , 6 , 0xA , 0xE , 2 , 4 , in18 , 0x80000000U );
128
+ blake32GS (3 , 7 , 0xB , 0xF , 10 , 15 , 0U , 640U );
129
+ blake32GS (0 , 5 , 0xA , 0xF , 14 , 1 , 0U , in17 );
130
+ blake32GS (1 , 6 , 0xB , 0xC , 11 , 12 , 0U , 0U );
131
+ blake32GS (2 , 7 , 0x8 , 0xD , 6 , 8 , 0U , 0U );
132
+ blake32GS (3 , 4 , 0x9 , 0xE , 3 , 13 , nonce , 1U );
133
+
134
+ // { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
135
+ blake32GS (0 , 4 , 0x8 , 0xC , 2 , 12 , in18 , 0U );
136
+ blake32GS (1 , 5 , 0x9 , 0xD , 6 , 10 , 0U , 0U );
137
+ blake32GS (2 , 6 , 0xA , 0xE , 0 , 11 , in16 , 0U );
138
+ blake32GS (3 , 7 , 0xB , 0xF , 8 , 3 , 0U , nonce );
139
+ blake32GS (0 , 5 , 0xA , 0xF , 4 , 13 , 0x80000000U , 1U );
140
+ blake32GS (1 , 6 , 0xB , 0xC , 7 , 5 , 0U , 0U );
141
+ blake32GS (2 , 7 , 0x8 , 0xD , 15 , 14 , 640U , 0U );
142
+ blake32GS (3 , 4 , 0x9 , 0xE , 1 , 9 , in17 , 0U );
143
+
144
+ // { 12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
145
+ blake32GS (0 , 4 , 0x8 , 0xC , 12 , 5 , 0U , 0U );
146
+ blake32GS (1 , 5 , 0x9 , 0xD , 1 , 15 , in17 , 640U );
147
+ blake32GS (2 , 6 , 0xA , 0xE , 14 , 13 , 0U , 1U );
148
+ blake32GS (3 , 7 , 0xB , 0xF , 4 , 10 , 0x80000000U , 0U );
149
+ blake32GS (0 , 5 , 0xA , 0xF , 0 , 7 , in16 , 0U );
150
+ blake32GS (1 , 6 , 0xB , 0xC , 6 , 3 , 0U , nonce );
151
+ blake32GS (2 , 7 , 0x8 , 0xD , 9 , 2 , 0U , in18 );
152
+ blake32GS (3 , 4 , 0x9 , 0xE , 8 , 11 , 0U , 0U );
153
+
154
+ // { 13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
155
+ blake32GS (0 , 4 , 0x8 , 0xC , 13 , 11 , 1U , 0U );
156
+ blake32GS (1 , 5 , 0x9 , 0xD , 7 , 14 , 0U , 0U );
157
+ blake32GS (2 , 6 , 0xA , 0xE , 12 , 1 , 0U , in17 );
158
+ blake32GS (3 , 7 , 0xB , 0xF , 3 , 9 , nonce , 0U );
159
+ blake32GS (0 , 5 , 0xA , 0xF , 5 , 0 , 0U , in16 );
160
+ blake32GS (1 , 6 , 0xB , 0xC , 15 , 4 , 640U , 0x80000000U );
161
+ blake32GS (2 , 7 , 0x8 , 0xD , 8 , 6 , 0U , 0U );
162
+ blake32GS (3 , 4 , 0x9 , 0xE , 2 , 10 , in18 , 0U );
163
+
164
+ // { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
165
+ blake32GS (0 , 4 , 0x8 , 0xC , 6 , 15 , 0U , 640U );
166
+ blake32GS (1 , 5 , 0x9 , 0xD , 14 , 9 , 0U , 0U );
167
+ blake32GS (2 , 6 , 0xA , 0xE , 11 , 3 , 0U , nonce );
168
+ blake32GS (3 , 7 , 0xB , 0xF , 0 , 8 , in16 , 0U );
169
+ blake32GS (0 , 5 , 0xA , 0xF , 12 , 2 , 0U , in18 );
170
+ blake32GS (1 , 6 , 0xB , 0xC , 13 , 7 , 1U , 0U );
171
+ blake32GS (2 , 7 , 0x8 , 0xD , 1 , 4 , in17 , 0x80000000U );
172
+ blake32GS (3 , 4 , 0x9 , 0xE , 10 , 5 , 0U , 0U );
173
+
174
+ // { 10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0 },
175
+ blake32GS (0 , 4 , 0x8 , 0xC , 10 , 2 , 0U , in18 );
176
+ blake32GS (1 , 5 , 0x9 , 0xD , 8 , 4 , 0U , 0x80000000U );
177
+ blake32GS (2 , 6 , 0xA , 0xE , 7 , 6 , 0U , 0U );
178
+ blake32GS (3 , 7 , 0xB , 0xF , 1 , 5 , in17 , 0U );
179
+ blake32GS (0 , 5 , 0xA , 0xF , 15 , 11 , 640U , 0U );
180
+ blake32GS (1 , 6 , 0xB , 0xC , 9 , 14 , 0U , 0U );
181
+ blake32GS (2 , 7 , 0x8 , 0xD , 3 , 12 , nonce , 0U );
182
+ blake32GS (3 , 4 , 0x9 , 0xE , 13 , 0 , 1U , in16 );
183
+
184
+
185
+ // { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
186
+ blake32GS (0 , 4 , 0x8 , 0xC , 0 , 1 , in16 , in17 );
187
+ blake32GS (1 , 5 , 0x9 , 0xD , 2 , 3 , in18 , nonce );
188
+ blake32GS (2 , 6 , 0xA , 0xE , 4 , 5 , 0x80000000U , 0U );
189
+ blake32GS (3 , 7 , 0xB , 0xF , 6 , 7 , 0U , 0U );
190
+ blake32GS (0 , 5 , 0xA , 0xF , 8 , 9 , 0U , 0U );
191
+ blake32GS (1 , 6 , 0xB , 0xC , 10 , 11 , 0U , 0U );
192
+ blake32GS (2 , 7 , 0x8 , 0xD , 12 , 13 , 0U , 1U );
193
+ blake32GS (3 , 4 , 0x9 , 0xE , 14 , 15 , 0U , 640U );
194
+
195
+ // { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
196
+ blake32GS (0 , 4 , 0x8 , 0xC , 14 , 10 , 0U , 0U );
197
+ blake32GS (1 , 5 , 0x9 , 0xD , 4 , 8 , 0x80000000 , 0U );
198
+ blake32GS (2 , 6 , 0xA , 0xE , 9 , 15 , 0U , 640U );
199
+ blake32GS (3 , 7 , 0xB , 0xF , 13 , 6 , 1U , 0U );
200
+ blake32GS (0 , 5 , 0xA , 0xF , 1 , 12 , in17 , 0U );
201
+ blake32GS (1 , 6 , 0xB , 0xC , 0 , 2 , in16 , in18 );
202
+ blake32GS (2 , 7 , 0x8 , 0xD , 11 , 7 , 0U , 0U );
203
+ blake32GS (3 , 4 , 0x9 , 0xE , 5 , 3 , 0U , nonce );
204
+
205
+ // { 11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
206
+ blake32GS (0 , 4 , 0x8 , 0xC , 11 , 8 , 0U , 0U );
207
+ blake32GS (1 , 5 , 0x9 , 0xD , 12 , 0 , 0U , in16 );
208
+ blake32GS (2 , 6 , 0xA , 0xE , 5 , 2 , 0U , in18 );
209
+ blake32GS (3 , 7 , 0xB , 0xF , 15 , 13 , 640U , 1U );
210
+ blake32GS (0 , 5 , 0xA , 0xF , 10 , 14 , 0U , 0U );
211
+ blake32GS (1 , 6 , 0xB , 0xC , 3 , 6 , nonce , 0U );
212
+ blake32GS (2 , 7 , 0x8 , 0xD , 7 , 1 , 0U , in17 );
213
+ blake32GS (3 , 4 , 0x9 , 0xE , 9 , 4 , 0U , 0x80000000U );
214
+
215
+ // { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
216
+ blake32GS (0 , 4 , 0x8 , 0xC , 7 , 9 , 0U , 0U );
217
+ blake32GS (1 , 5 , 0x9 , 0xD , 3 , 1 , nonce , in17 );
218
+ blake32GS (2 , 6 , 0xA , 0xE , 13 , 12 , 1U , 0U );
219
+ blake32GS (3 , 7 , 0xB , 0xF , 11 , 14 , 0U , 0U );
220
+ blake32GS (0 , 5 , 0xA , 0xF , 2 , 6 , in18 , 0U );
221
+ blake32GS (1 , 6 , 0xB , 0xC , 5 , 10 , 0U , 0U );
222
+ blake32GS (2 , 7 , 0x8 , 0xD , 4 , 0 , 0x80000000U , in16 );
223
+ blake32GS (3 , 4 , 0x9 , 0xE , 15 , 8 , 640U , 0U );
224
+
225
+
226
+ h [0 ] ^= v [0 ] ^ v [8 ];
227
+ h [1 ] ^= v [1 ] ^ v [9 ];
228
+ h [2 ] ^= v [2 ] ^ v [10 ];
229
+ h [3 ] ^= v [3 ] ^ v [11 ];
230
+ h [4 ] ^= v [4 ] ^ v [12 ];
231
+ h [5 ] ^= v [5 ] ^ v [13 ];
232
+ h [6 ] ^= v [6 ] ^ v [14 ];
233
+ h [7 ] ^= v [7 ] ^ v [15 ];
234
+
235
+ for (int i = 0 ; i < 8 ; ++ i )
236
+ {
237
+ byteSwapU32 (h [i ], h [i ]);
238
+ }
239
+
240
+ hash -> h4 [0 ] = (uint4 )(h [0 ], h [1 ], h [2 ], h [3 ]);
241
+ hash -> h4 [1 ] = (uint4 )(h [4 ], h [5 ], h [6 ], h [7 ]);
242
+
243
+ barrier (CLK_LOCAL_MEM_FENCE );
244
+ }
0 commit comments