PageRenderTime 51ms CodeModel.GetById 25ms RepoModel.GetById 0ms app.codeStats 0ms

/src/opencl/sha1_kernel.cl

https://bitbucket.org/lcirvin/cse465-project
OpenCL | 232 lines | 206 code | 26 blank | 0 comment | 0 complexity | da6c2833aa218ae5704ac038bc0596c4 MD5 | raw file
Possible License(s): BSD-3-Clause
  1. /*
  2. This code was largely inspired by
  3. pyrit opencl kernel sha1 routines, royger's sha1 sample,
  4. and md5_opencl_kernel.cl inside jtr.
  5. Copyright 2011 by Samuele Giovanni Tonon
  6. samu at linuxasylum dot net
  7. This program comes with ABSOLUTELY NO WARRANTY; express or
  8. implied .
  9. This is free software, and you are welcome to redistribute it
  10. under certain conditions; as expressed here
  11. http://www.gnu.org/licenses/gpl-2.0.html
  12. */
  13. #ifdef cl_khr_byte_addressable_store
  14. #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : disable
  15. #endif
  16. #ifdef cl_nv_pragma_unroll
  17. #define NVIDIA
  18. #pragma OPENCL EXTENSION cl_nv_pragma_unroll : enable
  19. #endif
  20. #ifdef NVIDIA
  21. inline uint SWAP32(uint x)
  22. {
  23. x = rotate(x, 16U);
  24. return ((x & 0x00FF00FF) << 8) + ((x >> 8) & 0x00FF00FF);
  25. }
  26. #else
  27. #define SWAP32(a) (as_uint(as_uchar4(a).wzyx))
  28. #endif
  29. #define K0 0x5A827999
  30. #define K1 0x6ED9EBA1
  31. #define K2 0x8F1BBCDC
  32. #define K3 0xCA62C1D6
  33. #define H1 0x67452301
  34. #define H2 0xEFCDAB89
  35. #define H3 0x98BADCFE
  36. #define H4 0x10325476
  37. #define H5 0xC3D2E1F0
  38. #ifndef uint32_t
  39. #define uint32_t unsigned int
  40. #endif
  41. __kernel void sha1_crypt_kernel(__global uint *data_info,__global char *plain_key, __global uint *digest){
  42. int t, gid, msg_pad;
  43. int stop, mmod;
  44. uint i, ulen;
  45. uint W[16], temp, A,B,C,D,E;
  46. uint num_keys = data_info[1];
  47. gid = get_global_id(0);
  48. msg_pad = gid * data_info[0];
  49. A = H1;
  50. B = H2;
  51. C = H3;
  52. D = H4;
  53. E = H5;
  54. #pragma unroll
  55. for (t = 1; t < 15; t++){
  56. W[t] = 0x00000000;
  57. }
  58. for(i = 0; i < data_info[0] && ((uchar) plain_key[msg_pad + i]) != 0x0 ; i++){
  59. }
  60. stop = i / 4 ;
  61. for (t = 0 ; t < stop ; t++){
  62. W[t] = ((uchar) plain_key[msg_pad + t * 4]) << 24;
  63. W[t] |= ((uchar) plain_key[msg_pad + t * 4 + 1]) << 16;
  64. W[t] |= ((uchar) plain_key[msg_pad + t * 4 + 2]) << 8;
  65. W[t] |= (uchar) plain_key[msg_pad + t * 4 + 3];
  66. }
  67. mmod = i % 4;
  68. if ( mmod == 3){
  69. W[t] = ((uchar) plain_key[msg_pad + t * 4]) << 24;
  70. W[t] |= ((uchar) plain_key[msg_pad + t * 4 + 1]) << 16;
  71. W[t] |= ((uchar) plain_key[msg_pad + t * 4 + 2]) << 8;
  72. W[t] |= ((uchar) 0x80) ;
  73. } else if (mmod == 2) {
  74. W[t] = ((uchar) plain_key[msg_pad + t * 4]) << 24;
  75. W[t] |= ((uchar) plain_key[msg_pad + t * 4 + 1]) << 16;
  76. W[t] |= 0x8000 ;
  77. } else if (mmod == 1) {
  78. W[t] = ((uchar) plain_key[msg_pad + t * 4]) << 24;
  79. W[t] |= 0x800000 ;
  80. } else /*if (mmod == 0)*/ {
  81. W[t] = 0x80000000 ;
  82. }
  83. ulen = (i * 8) & 0xFFFFFFFF;
  84. W[15] = ulen ;
  85. #undef R
  86. #define R(t) \
  87. ( \
  88. temp = W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \
  89. W[(t - 14) & 0x0F] ^ W[ t & 0x0F], \
  90. ( W[t & 0x0F] = rotate((int)temp,1) ) \
  91. )
  92. #undef P
  93. #define P(a,b,c,d,e,x) \
  94. { \
  95. e += rotate((int)a,5) + F(b,c,d) + K + x; b = rotate((int)b,30);\
  96. }
  97. #ifdef NVIDIA
  98. #define F(x,y,z) (z ^ (x & (y ^ z)))
  99. #else
  100. #define F(x,y,z) bitselect(z, y, x)
  101. #endif
  102. #define K 0x5A827999
  103. P( A, B, C, D, E, W[0] );
  104. P( E, A, B, C, D, W[1] );
  105. P( D, E, A, B, C, W[2] );
  106. P( C, D, E, A, B, W[3] );
  107. P( B, C, D, E, A, W[4] );
  108. P( A, B, C, D, E, W[5] );
  109. P( E, A, B, C, D, W[6] );
  110. P( D, E, A, B, C, W[7] );
  111. P( C, D, E, A, B, W[8] );
  112. P( B, C, D, E, A, W[9] );
  113. P( A, B, C, D, E, W[10] );
  114. P( E, A, B, C, D, W[11] );
  115. P( D, E, A, B, C, W[12] );
  116. P( C, D, E, A, B, W[13] );
  117. P( B, C, D, E, A, W[14] );
  118. P( A, B, C, D, E, W[15] );
  119. P( E, A, B, C, D, R(16) );
  120. P( D, E, A, B, C, R(17) );
  121. P( C, D, E, A, B, R(18) );
  122. P( B, C, D, E, A, R(19) );
  123. #undef K
  124. #undef F
  125. #define F(x,y,z) (x ^ y ^ z)
  126. #define K 0x6ED9EBA1
  127. P( A, B, C, D, E, R(20) );
  128. P( E, A, B, C, D, R(21) );
  129. P( D, E, A, B, C, R(22) );
  130. P( C, D, E, A, B, R(23) );
  131. P( B, C, D, E, A, R(24) );
  132. P( A, B, C, D, E, R(25) );
  133. P( E, A, B, C, D, R(26) );
  134. P( D, E, A, B, C, R(27) );
  135. P( C, D, E, A, B, R(28) );
  136. P( B, C, D, E, A, R(29) );
  137. P( A, B, C, D, E, R(30) );
  138. P( E, A, B, C, D, R(31) );
  139. P( D, E, A, B, C, R(32) );
  140. P( C, D, E, A, B, R(33) );
  141. P( B, C, D, E, A, R(34) );
  142. P( A, B, C, D, E, R(35) );
  143. P( E, A, B, C, D, R(36) );
  144. P( D, E, A, B, C, R(37) );
  145. P( C, D, E, A, B, R(38) );
  146. P( B, C, D, E, A, R(39) );
  147. #undef K
  148. #undef F
  149. #ifdef NVIDIA
  150. #define F(x,y,z) ((x & y) | (z & (x | y)))
  151. #else
  152. #define F(x,y,z) (bitselect(x, y, z) ^ bitselect(x, 0U, y))
  153. #endif
  154. #define K 0x8F1BBCDC
  155. P( A, B, C, D, E, R(40) );
  156. P( E, A, B, C, D, R(41) );
  157. P( D, E, A, B, C, R(42) );
  158. P( C, D, E, A, B, R(43) );
  159. P( B, C, D, E, A, R(44) );
  160. P( A, B, C, D, E, R(45) );
  161. P( E, A, B, C, D, R(46) );
  162. P( D, E, A, B, C, R(47) );
  163. P( C, D, E, A, B, R(48) );
  164. P( B, C, D, E, A, R(49) );
  165. P( A, B, C, D, E, R(50) );
  166. P( E, A, B, C, D, R(51) );
  167. P( D, E, A, B, C, R(52) );
  168. P( C, D, E, A, B, R(53) );
  169. P( B, C, D, E, A, R(54) );
  170. P( A, B, C, D, E, R(55) );
  171. P( E, A, B, C, D, R(56) );
  172. P( D, E, A, B, C, R(57) );
  173. P( C, D, E, A, B, R(58) );
  174. P( B, C, D, E, A, R(59) );
  175. #undef K
  176. #undef F
  177. #define F(x,y,z) (x ^ y ^ z)
  178. #define K 0xCA62C1D6
  179. P( A, B, C, D, E, R(60) );
  180. P( E, A, B, C, D, R(61) );
  181. P( D, E, A, B, C, R(62) );
  182. P( C, D, E, A, B, R(63) );
  183. P( B, C, D, E, A, R(64) );
  184. P( A, B, C, D, E, R(65) );
  185. P( E, A, B, C, D, R(66) );
  186. P( D, E, A, B, C, R(67) );
  187. P( C, D, E, A, B, R(68) );
  188. P( B, C, D, E, A, R(69) );
  189. P( A, B, C, D, E, R(70) );
  190. P( E, A, B, C, D, R(71) );
  191. P( D, E, A, B, C, R(72) );
  192. P( C, D, E, A, B, R(73) );
  193. P( B, C, D, E, A, R(74) );
  194. P( A, B, C, D, E, R(75) );
  195. P( E, A, B, C, D, R(76) );
  196. P( D, E, A, B, C, R(77) );
  197. P( C, D, E, A, B, R(78) );
  198. P( B, C, D, E, A, R(79) );
  199. #undef K
  200. #undef F
  201. digest[gid] = SWAP32(A + H1);
  202. digest[gid+1*num_keys] = SWAP32(B + H2);
  203. digest[gid+2*num_keys] = SWAP32(C + H3);
  204. digest[gid+3*num_keys] = SWAP32(D + H4);
  205. digest[gid+4*num_keys] = SWAP32(E + H5);
  206. }