diff --git a/tests/hash/wbag_gpu_bench.cu b/tests/hash/wbag_gpu_bench.cu index 55665be43..aa6f105fd 100644 --- a/tests/hash/wbag_gpu_bench.cu +++ b/tests/hash/wbag_gpu_bench.cu @@ -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<<>>(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 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;