#include <stdio.h>
#include <getopt.h>
#include <unistd.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <stdbool.h>
#include <sys/time.h>
typedef unsigned int uint;
typedef unsigned long long ull;
__device__ uint LR( uint a, int x) {
return a << x | a >> 32 - x;
}
__global__ void sha1_kernel( ull* res, ull IDX) {
ull id = threadIdx.x | ( ull) blockIdx.x << 10 | ( ull) blockIdx.y << 20 | ( ull) IDX << 32 , idx = id;
uint h0, h1, h2, h3, h4;
h0 = 0x63744d88 ; h1 = 0xf3eeebf1 ; h2 = 0x2e63d1c4 ; h3 = 0x761de6c4 ; h4 = 0x98de7acd ;
uint w[ 16 ] ;
for ( int i = 0 ; i < 8 ; i++ , idx >>= 16 ) {
w[ i] = 0x40404040u | ( idx& 15 ) << 24 | ( idx>> 4 & 15 ) << 16 | ( idx>> 8 & 15 ) << 8 | ( idx>> 12 & 15 ) ;
}
w[ 8 ] = 0x0a800000 ;
for ( int i = 9 ; i < 15 ; i++ ) w[ i] = 0 ;
w[ 15 ] = 1800 ;
uint a, b, c, d, e, f, k;
a = h0; b = h1; c = h2; d = h3; e = h4;
for ( int i = 0 ; i < 16 ; i++ ) {
f = ( b& c) | ( ~b& d) ;
k = 0x5A827999 ;
uint tmp = LR( a, 5 ) + f + e + k + w[ i& 15 ] ;
e = d; d = c; c = LR( b, 30 ) ; b = a; a = tmp;
}
for ( int i = 16 ; i < 20 ; i++ ) {
w[ i& 15 ] = LR( w[ i- 3 & 15 ] ^ w[ i- 8 & 15 ] ^ w[ i- 14 & 15 ] ^ w[ i& 15 ] , 1 ) ;
f = ( b& c) | ( ~b& d) ;
k = 0x5A827999 ;
uint tmp = LR( a, 5 ) + f + e + k + w[ i& 15 ] ;
e = d; d = c; c = LR( b, 30 ) ; b = a; a = tmp;
}
for ( int i = 20 ; i < 40 ; i++ ) {
w[ i& 15 ] = LR( w[ i- 3 & 15 ] ^ w[ i- 8 & 15 ] ^ w[ i- 14 & 15 ] ^ w[ i& 15 ] , 1 ) ;
f = ( b^ c^ d) ;
k = 0x6ED9EBA1 ;
uint tmp = LR( a, 5 ) + f + e + k + w[ i& 15 ] ;
e = d; d = c; c = LR( b, 30 ) ; b = a; a = tmp;
}
for ( int i = 40 ; i < 60 ; i++ ) {
w[ i& 15 ] = LR( w[ i- 3 & 15 ] ^ w[ i- 8 & 15 ] ^ w[ i- 14 & 15 ] ^ w[ i& 15 ] , 1 ) ;
f = ( b& c) | ( b& d) | ( c& d) ;
k = 0x8F1BBCDC ;
uint tmp = LR( a, 5 ) + f + e + k + w[ i& 15 ] ;
e = d; d = c; c = LR( b, 30 ) ; b = a; a = tmp;
}
for ( int i = 60 ; i < 80 ; i++ ) {
w[ i& 15 ] = LR( w[ i- 3 & 15 ] ^ w[ i- 8 & 15 ] ^ w[ i- 14 & 15 ] ^ w[ i& 15 ] , 1 ) ;
f = b^ c^ d;
k = 0xCA62C1D6 ;
uint tmp = LR( a, 5 ) + f + e + k + w[ i& 15 ] ;
e = d; d = c; c = LR( b, 30 ) ; b = a; a = tmp;
}
h0 = h0 + a;
h1 = h1 + b;
h2 = h2 + c;
h3 = h3 + d;
h4 = h4 + e;
if ( h0 == 0 && ( h1 >> 20 ) == 0 ) * res = id;
}
const int SZ = 4096 * 4096 ;
ull * res;
ull res_copy;
int main( )
{
dim3 threadsPerBlock( 1024 , 1 ) ;
dim3 numBlocks( 1024 , 4096 ) ;
cudaMalloc( & res, sizeof ( ull) ) ;
for ( int i = 200000 ;; i++ ) {
printf ( "IDX : %d\n " , i) ;
sha1_kernel<<< numBlocks, threadsPerBlock>>> ( res, i) ;
cudaMemcpy( & res_copy, res, sizeof ( ull) , cudaMemcpyDeviceToHost) ;
if ( res_copy == 0 ) continue ;
for ( int i = 0 ; i < 16 ; i++ ) {
printf ( "%c" , 64 | res_copy& 15 ) ;
res_copy / = 16 ;
}
break ;
}
}
I2luY2x1ZGUgPHN0ZGlvLmg+CiNpbmNsdWRlIDxnZXRvcHQuaD4KI2luY2x1ZGUgPHVuaXN0ZC5oPgojaW5jbHVkZSA8c3RkbGliLmg+CiNpbmNsdWRlIDxzdHJpbmcuaD4KI2luY2x1ZGUgPG1hdGguaD4KI2luY2x1ZGUgPHN0ZGJvb2wuaD4KI2luY2x1ZGUgPHN5cy90aW1lLmg+Cgp0eXBlZGVmIHVuc2lnbmVkIGludCB1aW50Owp0eXBlZGVmIHVuc2lnbmVkIGxvbmcgbG9uZyB1bGw7CgpfX2RldmljZV9fIHVpbnQgTFIodWludCBhLCBpbnQgeCl7CglyZXR1cm4gYSA8PCB4IHwgYSA+PiAzMi14Owp9CgpfX2dsb2JhbF9fIHZvaWQgc2hhMV9rZXJuZWwodWxsKiByZXMsIHVsbCBJRFgpIHsKCXVsbCBpZCA9IHRocmVhZElkeC54IHwgKHVsbClibG9ja0lkeC54IDw8IDEwIHwgKHVsbClibG9ja0lkeC55IDw8IDIwIHwgKHVsbClJRFggPDwgMzIsIGlkeCA9IGlkOwoJdWludCBoMCwgaDEsIGgyLCBoMywgaDQ7CgloMCA9IDB4NjM3NDRkODg7IGgxID0gMHhmM2VlZWJmMTsgaDIgPSAweDJlNjNkMWM0OyBoMyA9IDB4NzYxZGU2YzQ7IGg0ID0gMHg5OGRlN2FjZDsKCgl1aW50IHdbMTZdOwoJZm9yKGludCBpID0gMDsgaSA8IDg7IGkrKywgaWR4ID4+PSAxNil7CgkJd1tpXSA9IDB4NDA0MDQwNDB1IHwgKGlkeCYxNSkgPDwgMjQgfCAoaWR4Pj40JjE1KSA8PCAxNiB8IChpZHg+PjgmMTUpIDw8IDggfCAoaWR4Pj4xMiYxNSk7Cgl9Cgl3WzhdID0gMHgwYTgwMDAwMDsKCWZvcihpbnQgaSA9IDk7IGkgPCAxNTsgaSsrKSB3W2ldID0gMDsKCXdbMTVdID0gMTgwMDsKCgl1aW50IGEsIGIsIGMsIGQsIGUsIGYsIGs7CglhID0gaDA7IGIgPSBoMTsgYyA9IGgyOyBkID0gaDM7IGUgPSBoNDsKCWZvcihpbnQgaSA9IDA7IGkgPCAxNjsgaSsrKXsKCQlmID0gKGImYyl8KH5iJmQpOwoJCWsgPSAweDVBODI3OTk5OwoJCXVpbnQgdG1wID0gTFIoYSwgNSkgKyBmICsgZSArIGsgKyB3W2kmMTVdOwoJCWUgPSBkOyBkID0gYzsgYyA9IExSKGIsIDMwKTsgYiA9IGE7IGEgPSB0bXA7Cgl9Cglmb3IoaW50IGkgPSAxNjsgaSA8IDIwOyBpKyspewoJCXdbaSYxNV0gPSBMUih3W2ktMyYxNV1ed1tpLTgmMTVdXndbaS0xNCYxNV1ed1tpJjE1XSwgMSk7CgoJCWYgPSAoYiZjKXwofmImZCk7CgkJayA9IDB4NUE4Mjc5OTk7CgoJCXVpbnQgdG1wID0gTFIoYSwgNSkgKyBmICsgZSArIGsgKyB3W2kmMTVdOwoJCWUgPSBkOyBkID0gYzsgYyA9IExSKGIsIDMwKTsgYiA9IGE7IGEgPSB0bXA7Cgl9Cglmb3IoaW50IGkgPSAyMDsgaSA8IDQwOyBpKyspewoJCXdbaSYxNV0gPSBMUih3W2ktMyYxNV1ed1tpLTgmMTVdXndbaS0xNCYxNV1ed1tpJjE1XSwgMSk7CgoJCWYgPSAoYl5jXmQpOwoJCWsgPSAweDZFRDlFQkExOyAKCgkJdWludCB0bXAgPSBMUihhLCA1KSArIGYgKyBlICsgayArIHdbaSYxNV07CgkJZSA9IGQ7IGQgPSBjOyBjID0gTFIoYiwgMzApOyBiID0gYTsgYSA9IHRtcDsKCX0KCWZvcihpbnQgaSA9IDQwOyBpIDwgNjA7IGkrKyl7CgkJd1tpJjE1XSA9IExSKHdbaS0zJjE1XV53W2ktOCYxNV1ed1tpLTE0JjE1XV53W2kmMTVdLCAxKTsKCgkJZiA9IChiJmMpfChiJmQpfChjJmQpOwoJCWsgPSAweDhGMUJCQ0RDOwoKCQl1aW50IHRtcCA9IExSKGEsIDUpICsgZiArIGUgKyBrICsgd1tpJjE1XTsKCQllID0gZDsgZCA9IGM7IGMgPSBMUihiLCAzMCk7IGIgPSBhOyBhID0gdG1wOwoJfQoJZm9yKGludCBpID0gNjA7IGkgPCA4MDsgaSsrKXsKCQl3W2kmMTVdID0gTFIod1tpLTMmMTVdXndbaS04JjE1XV53W2ktMTQmMTVdXndbaSYxNV0sIDEpOwoKCQlmID0gYl5jXmQ7CgkJayA9IDB4Q0E2MkMxRDY7CgoJCXVpbnQgdG1wID0gTFIoYSwgNSkgKyBmICsgZSArIGsgKyB3W2kmMTVdOwoJCWUgPSBkOyBkID0gYzsgYyA9IExSKGIsIDMwKTsgYiA9IGE7IGEgPSB0bXA7Cgl9CgloMCA9IGgwICsgYTsKCWgxID0gaDEgKyBiOwoJaDIgPSBoMiArIGM7CgloMyA9IGgzICsgZDsKCWg0ID0gaDQgKyBlOwoJaWYoaDAgPT0gMCAmJiAoaDEgPj4gMjApID09IDApICpyZXMgPSBpZDsKfQoKY29uc3QgaW50IFNaID0gNDA5NiAqIDQwOTY7Cgp1bGwgKnJlczsKdWxsIHJlc19jb3B5OwoKaW50IG1haW4oKQp7CglkaW0zIHRocmVhZHNQZXJCbG9jaygxMDI0LCAxKTsKCWRpbTMgbnVtQmxvY2tzKDEwMjQsIDQwOTYpOwoKCWN1ZGFNYWxsb2MoJnJlcywgc2l6ZW9mKHVsbCkpOwoKCWZvcihpbnQgaSA9IDIwMDAwMDs7IGkrKyl7CgkJcHJpbnRmKCJJRFggOiAlZFxuIiwgaSk7CgkJc2hhMV9rZXJuZWw8PDxudW1CbG9ja3MsIHRocmVhZHNQZXJCbG9jaz4+PihyZXMsIGkpOwoKCQljdWRhTWVtY3B5KCZyZXNfY29weSwgcmVzLCBzaXplb2YodWxsKSwgY3VkYU1lbWNweURldmljZVRvSG9zdCk7CgkJaWYocmVzX2NvcHkgPT0gMCkgY29udGludWU7CgkJZm9yKGludCBpID0gMDsgaSA8IDE2OyBpKyspewoJCQlwcmludGYoIiVjIiwgNjQgfCByZXNfY29weSYxNSk7CgkJCXJlc19jb3B5IC89IDE2OwoJCX0KCQlicmVhazsKCX0KfQo=
compilation info
prog.cpp:13:1: error: ‘__device__’ does not name a type
__device__ uint LR(uint a, int x){
^~~~~~~~~~
prog.cpp:17:1: error: ‘__global__’ does not name a type
__global__ void sha1_kernel(ull* res, ull IDX) {
^~~~~~~~~~
prog.cpp: In function ‘int main()’:
prog.cpp:89:2: error: ‘dim3’ was not declared in this scope
dim3 threadsPerBlock(1024, 1);
^~~~
prog.cpp:90:7: error: expected ‘;’ before ‘numBlocks’
dim3 numBlocks(1024, 4096);
^~~~~~~~~
prog.cpp:92:30: error: ‘cudaMalloc’ was not declared in this scope
cudaMalloc(&res, sizeof(ull));
^
prog.cpp:96:3: error: ‘sha1_kernel’ was not declared in this scope
sha1_kernel<<<numBlocks, threadsPerBlock>>>(res, i);
^~~~~~~~~~~
prog.cpp:96:16: error: expected primary-expression before ‘<’ token
sha1_kernel<<<numBlocks, threadsPerBlock>>>(res, i);
^
prog.cpp:96:17: error: ‘numBlocks’ was not declared in this scope
sha1_kernel<<<numBlocks, threadsPerBlock>>>(res, i);
^~~~~~~~~
prog.cpp:96:28: error: ‘threadsPerBlock’ was not declared in this scope
sha1_kernel<<<numBlocks, threadsPerBlock>>>(res, i);
^~~~~~~~~~~~~~~
prog.cpp:96:45: error: expected primary-expression before ‘>’ token
sha1_kernel<<<numBlocks, threadsPerBlock>>>(res, i);
^
prog.cpp:98:43: error: ‘cudaMemcpyDeviceToHost’ was not declared in this scope
cudaMemcpy(&res_copy, res, sizeof(ull), cudaMemcpyDeviceToHost);
^~~~~~~~~~~~~~~~~~~~~~
prog.cpp:98:65: error: ‘cudaMemcpy’ was not declared in this scope
cudaMemcpy(&res_copy, res, sizeof(ull), cudaMemcpyDeviceToHost);
^
stdout