Skip to content

Commit

Permalink
v 0.4.1, printing wif, stride in const memory
Browse files Browse the repository at this point in the history
  • Loading branch information
PawelGorny committed Jan 23, 2022
1 parent 8795f10 commit d813434
Show file tree
Hide file tree
Showing 5 changed files with 47 additions and 32 deletions.
12 changes: 7 additions & 5 deletions WifSolverCuda/Worker.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,10 +7,10 @@
#include "lib/Math.cuh"


__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks);
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks);
__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks, const uint32_t checksum);
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks, const uint32_t checksum);
__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks);
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks);
__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks, const uint32_t checksum);
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks, const uint32_t checksum);
__global__ void resultCollector(bool* buffResult, uint64_t* buffCombinedResult, const uint64_t threadsInBlockNumberOfChecks);

__device__ bool _checksumDoubleSha256CheckUncompressed(unsigned int checksum, beu32* d_hash, uint64_t* _start);
Expand All @@ -21,4 +21,6 @@ __device__ bool _checksumDoubleSha256(unsigned int checksum, beu32* d_hash);
__device__ void _add(uint64_t* C, uint64_t* A);
__device__ void _load(uint64_t* C, uint64_t* A);

__device__ void IMult(uint64_t* r, uint64_t* a, int64_t b);
__device__ void IMult(uint64_t* r, uint64_t* a, int64_t b);

cudaError_t loadStride(uint64_t* stride);
21 changes: 9 additions & 12 deletions WifSolverCuda/Worker1.cu
View file
Original file line number Diff line number Diff line change
@@ -1,12 +1,12 @@
#include "Worker.cuh"

__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks, const uint32_t checksum) {
uint64_t _stride[5];
__device__ __constant__ uint64_t _stride[5];

__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks, const uint32_t checksum) {
uint64_t _start[5];
uint64_t _startStride[5];
beu32 d_hash[8];
_load(_start, buffRangeStart);
_load(_stride, buffStride);

int64_t tIx = (threadIdx.x + blockIdx.x * blockDim.x) * threadNumberOfChecks;
IMult(_startStride, _stride, tIx);
Expand All @@ -19,13 +19,11 @@ __global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, ui
_add(_start, _stride);
}
}
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks, const uint32_t checksum) {
uint64_t _stride[5];
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks, const uint32_t checksum) {
uint64_t _start[5];
uint64_t _startStride[5];
beu32 d_hash[8];
_load(_start, buffRangeStart);
_load(_stride, buffStride);

int64_t tIx = (threadIdx.x + blockIdx.x * blockDim.x) * threadNumberOfChecks;
IMult(_startStride, _stride, tIx);
Expand All @@ -42,13 +40,11 @@ __global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint
_add(_start, _stride);
}
}
__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks) {
uint64_t _stride[5];
__global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks) {
uint64_t _start[5];
uint64_t _startStride[5];
beu32 d_hash[8];
_load(_start, buffRangeStart);
_load(_stride, buffStride);

int64_t tIx = (threadIdx.x + blockIdx.x * blockDim.x) * threadNumberOfChecks;
IMult(_startStride, _stride, tIx);
Expand All @@ -61,13 +57,11 @@ __global__ void kernelUncompressed(bool* buffResult, bool* buffCollectorWork, ui
_add(_start, _stride);
}
}
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* buffRangeStart, uint64_t* buffStride, const int threadNumberOfChecks) {
uint64_t _stride[5];
__global__ void kernelCompressed(bool* buffResult, bool* buffCollectorWork, uint64_t* const __restrict__ buffRangeStart, const int threadNumberOfChecks) {
uint64_t _start[5];
uint64_t _startStride[5];
beu32 d_hash[8];
_load(_start, buffRangeStart);
_load(_stride, buffStride);

int64_t tIx = (threadIdx.x + blockIdx.x * blockDim.x) * threadNumberOfChecks;
IMult(_startStride, _stride, tIx);
Expand Down Expand Up @@ -194,3 +188,6 @@ __device__ void IMult(uint64_t* r, uint64_t* a, int64_t b) {
Mult2(r, t, b)
}

cudaError_t loadStride(uint64_t* stride){
return cudaMemcpyToSymbol(_stride, stride, 5 * sizeof(uint64_t));
}
6 changes: 5 additions & 1 deletion WifSolverCuda/lib/util.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -189,11 +189,15 @@ void addressToBase58(char* rmd, char* dst) {
memcpy(digest + 1, rmd, 20);
sha256((uint8_t*)digest, 21, (uint8_t*)digest + 21);
sha256((uint8_t*)digest + 21, 32, (uint8_t*)digest + 21);
if (!b58enc(dst, &pubaddress_size, digest, 25)) {
if (!b58encode(dst, &pubaddress_size, digest, 25)) {
fprintf(stderr, "error b58enc\n");
}
}

bool b58encode(char* b58, size_t* b58sz, const void* data, size_t binsz) {
return b58enc(b58, b58sz, data, binsz);
}

std::string formatDouble(const char* formatStr, double value)
{
char buf[100] = { 0 };
Expand Down
2 changes: 2 additions & 0 deletions WifSolverCuda/lib/util.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -36,4 +36,6 @@ void addressToBase58(char* rmd, char* dst);

std::string formatDouble(const char* formatStr, double value);

bool b58encode(char* b58, size_t* b58sz, const void* data, size_t binsz);

#endif // CUSTOMUTILH
38 changes: 24 additions & 14 deletions WifSolverCuda/main.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,7 @@ Secp256K1* secp;

int main(int argc, char** argv)
{
printf("WifSolver 0.4.0\n\n");
printf("WifSolver 0.4.1\n\n");

if (readArgs(argc, argv)) {
showHelp();
Expand Down Expand Up @@ -98,22 +98,19 @@ cudaError_t processCuda() {
uint64_t* buffRangeStart = new uint64_t[NB64BLOCK];
uint64_t* dev_buffRangeStart = new uint64_t[NB64BLOCK];
uint64_t* buffStride = new uint64_t[NB64BLOCK];
uint64_t* dev_buffStride = new uint64_t[NB64BLOCK];

int COLLECTOR_SIZE = BLOCK_NUMBER;

__Load(buffStride, STRIDE.bits64);
loadStride(buffStride);

bool* buffDeviceResult = new bool[outputSize];
bool* dev_buffDeviceResult = new bool[outputSize];
for (int i = 0; i < outputSize; i++) {
buffDeviceResult[i] = false;
}
cudaStatus = cudaMalloc((void**)&dev_buffDeviceResult, outputSize * sizeof(bool));
cudaStatus = cudaMemcpyAsync(dev_buffDeviceResult, buffDeviceResult, outputSize * sizeof(bool), cudaMemcpyHostToDevice);

cudaStatus = cudaMalloc((void**)&dev_buffStride, NB64BLOCK * sizeof(uint64_t));
cudaStatus = cudaMemcpy(dev_buffStride, buffStride, NB64BLOCK * sizeof(uint64_t), cudaMemcpyHostToDevice);
cudaStatus = cudaMemcpyAsync(dev_buffDeviceResult, buffDeviceResult, outputSize * sizeof(bool), cudaMemcpyHostToDevice);

uint64_t* buffResult = new uint64_t[COLLECTOR_SIZE];
uint64_t* dev_buffResult = new uint64_t[COLLECTOR_SIZE];
Expand Down Expand Up @@ -143,16 +140,16 @@ cudaError_t processCuda() {
//launch work
if (COMPRESSED) {
if (IS_CHECKSUM) {
kernelCompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, dev_buffStride, THREAD_STEPS, expectedChecksum);
kernelCompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, THREAD_STEPS, expectedChecksum);
}else{
kernelCompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, dev_buffStride, THREAD_STEPS);
kernelCompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, THREAD_STEPS);
}
}
else {
if (IS_CHECKSUM) {
kernelUncompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, dev_buffStride, THREAD_STEPS, expectedChecksum);
kernelUncompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, THREAD_STEPS, expectedChecksum);
}else{
kernelUncompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, dev_buffStride, THREAD_STEPS);
kernelUncompressed << <BLOCK_NUMBER, BLOCK_THREADS >> > (dev_buffDeviceResult, dev_buffCollectorWork, dev_buffRangeStart, THREAD_STEPS);
}

}
Expand Down Expand Up @@ -255,15 +252,20 @@ Error:
cudaFree(dev_buffResult);
cudaFree(dev_buffDeviceResult);
cudaFree(dev_buffRangeStart);
cudaFree(dev_buffStride);
cudaFree(dev_buffCollectorWork);
return cudaStatus;
}


void processCandidate(Int &toTest) {
FILE* keys;
char rmdhash[21], address[50];
size_t dataLen = COMPRESSED ? 38 : 37;
size_t wifLen = 53;
char rmdhash[21], address[50], wif[53];
unsigned char* buff = new unsigned char[dataLen];
for (int i = 0, d=dataLen-1; i < dataLen; i++, d--) {
buff[i] = toTest.GetByte(d);
}
toTest.SetBase16((char*)toTest.GetBase16().substr(2, 64).c_str());
Point publickey = secp->ComputePublicKey(&toTest);
secp->GetHash160(P2PKH, COMPRESSED, publickey, (unsigned char*)rmdhash);
Expand All @@ -274,9 +276,13 @@ void processCandidate(Int &toTest) {
printf("\n");
printf("found: %s\n", address);
printf("key : %s\n", toTest.GetBase16().c_str());
if (b58encode(wif, &wifLen, buff, dataLen)) {
printf("WIF : %s\n", wif);
}
keys = fopen(fileResult.c_str(), "a+");
fprintf(keys, "%s\n", address);
fprintf(keys, "%s\n", toTest.GetBase16().c_str());
fprintf(keys, "%s\n", wif);
fprintf(keys, "%s\n\n", toTest.GetBase16().c_str());
fclose(keys);
return;
}
Expand All @@ -285,9 +291,13 @@ void processCandidate(Int &toTest) {
printf("\n");
printf("found: %s\n", address);
printf("key : %s\n", toTest.GetBase16().c_str());
if (b58encode(wif, &wifLen, buff, dataLen)) {
printf("WIF : %s\n", wif);
}
keys = fopen(fileResultPartial.c_str(), "a+");
fprintf(keys, "%s\n", address);
fprintf(keys, "%s\n", toTest.GetBase16().c_str());
fprintf(keys, "%s\n", wif);
fprintf(keys, "%s\n\n", toTest.GetBase16().c_str());
fclose(keys);
}
}
Expand Down

0 comments on commit d813434

Please sign in to comment.