codeberg-wowario/wownero
1
0
Fork 0
mirror of https://codeberg.org/wowario/wownero synced 2026-11-11 08:44:15 +01:00
Code Issues Packages Projects Releases Wiki Activity

update gpu code

Browse Source
This commit is contained in:
wowario
2025-09-26 11:24:00 +03:00
parent 311504607e
commit 65e649255b
1 changed files with 27 additions and 103 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/hash/wbag_gpu_bench.cu
+27 -103
View File
@@ -10,48 +10,24 @@
#include "crypto/wbag-hash.h"
#include "cryptonote_basic/wbag_cache.h"
__device__ __forceinline__ uint64_t d_rotl64(uint64_t x, unsigned r) {
return (x << (r & 63)) | (x >> ((64 - r) & 63));
}
__device__ __forceinline__ uint64_t d_splitmix64(uint64_t x) {
x += 0x9e3779b97f4a7c15ULL;
x = (x ^ (x >> 30)) * 0xbf58476d1ce4e5b9ULL;
x = (x ^ (x >> 27)) * 0x94d049bb133111ebULL;
return x ^ (x >> 31);
}
__device__ __forceinline__ uint64_t d_mulhi64(uint64_t a, uint64_t b) {
return __umul64hi(a, b);
}
__device__ __forceinline__ uint64_t d_load64_le(const uint8_t* p) {
uint64_t v; memcpy(&v, p, 8); return v;
}
__device__ __forceinline__ void d_store64_le(uint8_t* p, uint64_t x) {
memcpy(p, &x, 8);
}
__device__ __forceinline__ uint64_t d_rotl64(uint64_t x, unsigned r){ return (x << (r & 63)) | (x >> ((64 - r) & 63)); }
__device__ __forceinline__ uint64_t d_splitmix64(uint64_t x){ x += 0x9e3779b97f4a7c15ULL; x = (x ^ (x >> 30)) * 0xbf58476d1ce4e5b9ULL; x = (x ^ (x >> 27)) * 0x94d049bb133111ebULL; return x ^ (x >> 31); }
__device__ __forceinline__ uint64_t d_mulhi64(uint64_t a, uint64_t b){ return __umul64hi(a,b); }
__device__ __forceinline__ uint64_t d_load64_le(const uint8_t* p){ uint64_t v; memcpy(&v, p, 8); return v; }
__device__ __forceinline__ void d_store64_le(uint8_t* p, uint64_t x){ memcpy(p, &x, 8); }
__global__ void wbag_fill_resident(uint8_t* resident,
size_t resident_len,
uint64_t seed,
int nthreads,
int per_thread)
{
__global__ void wbag_fill_resident(uint8_t* resident, size_t resident_len, uint64_t seed, int nthreads, int per_thread){
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (per_thread) {
if (tid >= nthreads) return;
uint8_t* base = resident + (size_t)tid * resident_len;
uint64_t q = seed ^ 0x243F6A8885A308D3ULL;
for (size_t i = 0; i < resident_len; i += 8) {
q = d_splitmix64(q);
d_store64_le(base + i, q);
}
for (size_t i = 0; i < resident_len; i += 8){ q = d_splitmix64(q); d_store64_le(base + i, q); }
} else {
if (tid != 0) return;
uint8_t* base = resident;
uint64_t q = seed ^ 0x243F6A8885A308D3ULL;
for (size_t i = 0; i < resident_len; i += 8) {
q = d_splitmix64(q);
d_store64_le(base + i, q);
}
for (size_t i = 0; i < resident_len; i += 8){ q = d_splitmix64(q); d_store64_le(base + i, q); }
}
}
@@ -67,19 +43,13 @@ __global__ void wbag_kernel(const uint8_t* __restrict__ cache,
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= nthreads) return;
const uint8_t* my_resident = per_thread ? resident + (size_t)tid * resident_len
: resident;
const size_t n_mask = n_entries - 1;
const uint8_t* my_resident = per_thread ? resident + (size_t)tid * resident_len : resident;
const size_t n_mask = n_entries - 1;
const size_t smask_touch = (scratchpad_len / 8) - 1;
const size_t window_span = (resident_len > scratchpad_len)
? (resident_len - scratchpad_len) : 0;
uint64_t s = d_splitmix64(base_seed ^ (uint64_t)tid);
uint64_t st0 = d_splitmix64(s);
uint64_t st1 = d_splitmix64(st0);
uint64_t st2 = d_splitmix64(st1);
uint64_t st3 = d_splitmix64(st2);
const size_t window_span = (resident_len > scratchpad_len) ? (resident_len - scratchpad_len) : 0;
uint64_t s = d_splitmix64(base_seed ^ (uint64_t)tid);
uint64_t st0 = d_splitmix64(s), st1 = d_splitmix64(st0), st2 = d_splitmix64(st1), st3 = d_splitmix64(st2);
uint64_t checksum = 0;
for (int ttry = 0; ttry < trials; ++ttry) {
@@ -124,7 +94,7 @@ __global__ void wbag_kernel(const uint8_t* __restrict__ cache,
size_t window_base = 0;
if (window_span) {
const size_t align = window_align;
const size_t align = window_align;
const size_t max_base = window_span & ~(align - 1u);
const uint64_t rnd = d_rotl64(ss ^ mh_mix, 17) ^ d_rotl64(ss + (uint64_t)t, 41);
window_base = (size_t)(rnd % (max_base + 1u));
@@ -160,24 +130,16 @@ static inline int parse_resident_mode_env() {
#endif
}
int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view,
int threads, int trials, double& hashes_per_sec_out)
int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view, int threads, int trials, double& hashes_per_sec_out)
{
if (threads <= 0 || trials <= 0) return 1;
cudaError_t st = cudaFree(0);
if (st != cudaSuccess) {
std::fprintf(stderr, "CUDA init: %s\n", cudaGetErrorString(st));
return 100;
}
int devs = 0;
st = cudaGetDeviceCount(&devs);
cudaError_t st = cudaGetDeviceCount(&devs);
if (st != cudaSuccess || devs <= 0) {
std::fprintf(stderr, "CUDA get device count: %s\n", cudaGetErrorString(st));
return 100;
}
st = cudaSetDevice(0);
if (st != cudaSuccess) {
std::fprintf(stderr, "CUDA set device: %s\n", cudaGetErrorString(st));
@@ -199,8 +161,7 @@ int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view,
if (d_cache) { cudaFree(d_cache); d_cache = nullptr; d_cache_size = 0; }
st = cudaMalloc(&d_cache, cache_bytes);
if (st != cudaSuccess) {
std::fprintf(stderr, "GPU: cache alloc failed (%zu bytes): %s\n",
cache_bytes, cudaGetErrorString(st));
std::fprintf(stderr, "GPU: cache alloc failed (%zu bytes): %s\n", cache_bytes, cudaGetErrorString(st));
return 2;
}
st = cudaMemcpy(d_cache, view.data(), cache_bytes, cudaMemcpyHostToDevice);
@@ -222,18 +183,14 @@ int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view,
return 4;
}
if (need_resident + (64ULL << 20) > free_b) {
std::fprintf(stderr,
"GPU: not enough VRAM for resident (%zu MiB needed, %zu MiB free). "
"Set WBAG_GPU_RESIDENT=shared or reduce --threads.\n",
need_resident >> 20, free_b >> 20);
std::fprintf(stderr, "GPU: not enough VRAM for resident (%zu MiB needed, %zu MiB free). Set WBAG_GPU_RESIDENT=shared or reduce --threads.\n", need_resident >> 20, free_b >> 20);
return 5;
}
uint8_t* d_resident = nullptr;
st = cudaMalloc(&d_resident, need_resident);
if (st != cudaSuccess) {
std::fprintf(stderr, "GPU: resident alloc failed (%zu bytes): %s\n",
need_resident, cudaGetErrorString(st));
std::fprintf(stderr, "GPU: resident alloc failed (%zu bytes): %s\n", need_resident, cudaGetErrorString(st));
return 6;
}
@@ -244,26 +201,14 @@ int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view,
const int grid = per_thread ? ((threads + block - 1) / block) : 1;
wbag_fill_resident<<<grid, block>>>(d_resident, resident_len, base_seed, threads, per_thread);
cudaError_t e1 = cudaGetLastError();
if (e1 != cudaSuccess) {
std::fprintf(stderr, "GPU: fill launch failed: %s\n", cudaGetErrorString(e1));
cudaFree(d_resident);
return 7;
}
if (e1 != cudaSuccess) { std::fprintf(stderr, "GPU: fill launch failed: %s\n", cudaGetErrorString(e1)); cudaFree(d_resident); return 7; }
e1 = cudaDeviceSynchronize();
if (e1 != cudaSuccess) {
std::fprintf(stderr, "GPU: fill runtime error: %s\n", cudaGetErrorString(e1));
cudaFree(d_resident);
return 8;
}
if (e1 != cudaSuccess) { std::fprintf(stderr, "GPU: fill runtime error: %s\n", cudaGetErrorString(e1)); cudaFree(d_resident); return 8; }
}
uint64_t* d_out = nullptr;
st = cudaMalloc(&d_out, sizeof(uint64_t) * (size_t)threads);
if (st != cudaSuccess) {
std::fprintf(stderr, "GPU: out alloc failed: %s\n", cudaGetErrorString(st));
cudaFree(d_resident);
return 9;
}
if (st != cudaSuccess) { std::fprintf(stderr, "GPU: out alloc failed: %s\n", cudaGetErrorString(st)); cudaFree(d_resident); return 9; }
cudaEvent_t ev0, ev1; cudaEventCreate(&ev0); cudaEventCreate(&ev1);
cudaEventRecord(ev0);
@@ -277,39 +222,18 @@ int wbag_gpu_bench_once(const cryptonote::wbag::CacheView& view,
base_seed, trials, threads, d_out, per_thread);
cudaError_t e1 = cudaGetLastError();
if (e1 != cudaSuccess) {
std::fprintf(stderr, "GPU: hash launch failed: %s\n", cudaGetErrorString(e1));
cudaFree(d_out); cudaFree(d_resident);
cudaEventDestroy(ev0); cudaEventDestroy(ev1);
return 10;
}
if (e1 != cudaSuccess) { std::fprintf(stderr, "GPU: hash launch failed: %s\n", cudaGetErrorString(e1)); cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); return 10; }
e1 = cudaDeviceSynchronize();
if (e1 != cudaSuccess) {
std::fprintf(stderr, "GPU: hash runtime error: %s\n", cudaGetErrorString(e1));
cudaFree(d_out); cudaFree(d_resident);
cudaEventDestroy(ev0); cudaEventDestroy(ev1);
return 11;
}
if (e1 != cudaSuccess) { std::fprintf(stderr, "GPU: hash runtime error: %s\n", cudaGetErrorString(e1)); cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); return 11; }
cudaEventRecord(ev1);
cudaEventSynchronize(ev1);
cudaEventRecord(ev1); cudaEventSynchronize(ev1);
float ms = 0.f; cudaEventElapsedTime(&ms, ev0, ev1);
std::vector<uint64_t> h_out((size_t)threads);
st = cudaMemcpy(h_out.data(), d_out, sizeof(uint64_t) * (size_t)threads, cudaMemcpyDeviceToHost);
if (st != cudaSuccess) {
std::fprintf(stderr, "GPU: out memcpy failed: %s\n", cudaGetErrorString(st));
cudaFree(d_out); cudaFree(d_resident);
cudaEventDestroy(ev0); cudaEventDestroy(ev1);
return 12;
}
if (st != cudaSuccess) { std::fprintf(stderr, "GPU: out memcpy failed: %s\n", cudaGetErrorString(st)); cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); return 12; }
uint64_t acc = 0; for (uint64_t v : h_out) acc |= v;
if (!acc) {
std::fprintf(stderr, "GPU: zero checksum\n");
cudaFree(d_out); cudaFree(d_resident);
cudaEventDestroy(ev0); cudaEventDestroy(ev1);
return 13;
}
if (!acc) { std::fprintf(stderr, "GPU: zero checksum\n"); cudaFree(d_out); cudaFree(d_resident); cudaEventDestroy(ev0); cudaEventDestroy(ev1); return 13; }
const double sec = std::max(1e-6, double(ms) / 1000.0);
hashes_per_sec_out = (double(trials) * double(threads)) / sec;