// // =============== BLAKE part on nVidia GPU ====================== // // This is the generic "default" implementation when no architecture // specific implementation is available in the kernel. // // NOTE: compile this .cu module for compute_10,sm_10 with --maxrregcount=64 // // TODO: CUDA porting work remains to be done. // #include <map> #include <stdint.h> #include "miner.h" #include "salsa_kernel.h" #include "cuda_helper.h" typedef uint32_t sph_u32; #define SPH_ROTL32 ROTL32 #define SPH_ROTR32 ROTR32 __constant__ uint64_t ptarget64[4]; __constant__ uint32_t pdata[20]; // define some error checking macros #define DELIMITER '/' #define __FILENAME__ ( strrchr(__FILE__, DELIMITER) != NULL ? strrchr(__FILE__, DELIMITER)+1 : __FILE__ ) #undef checkCudaErrors #define checkCudaErrors(x) \ { \ cudaGetLastError(); \ x; \ cudaError_t err = cudaGetLastError(); \ if (err != cudaSuccess && !abort_flag) \ applog(LOG_ERR, "GPU #%d: cudaError %d (%s) (%s line %d)\n", device_map[thr_id], err, cudaGetErrorString(err), __FILENAME__, __LINE__); \ } // from salsa_kernel.cu extern std::map<int, uint32_t *> context_idata[2]; extern std::map<int, uint32_t *> context_odata[2]; extern std::map<int, cudaStream_t> context_streams[2]; extern std::map<int, uint32_t *> context_hash[2]; #ifdef _MSC_VER #pragma warning (disable: 4146) #endif /** * Encode a 32-bit value into the provided buffer (big endian convention). * * @param dst the destination buffer * @param val the 32-bit value to encode */ static __device__ void cuda_sph_enc32be(void *dst, sph_u32 val) { *(sph_u32 *)dst = cuda_swab32(val); } #define Z00 0 #define Z01 1 #define Z02 2 #define Z03 3 #define Z04 4 #define Z05 5 #define Z06 6 #define Z07 7 #define Z08 8 #define Z09 9 #define Z0A A #define Z0B B #define Z0C C #define Z0D D #define Z0E E #define Z0F F #define Z10 E #define Z11 A #define Z12 4 #define Z13 8 #define Z14 9 #define Z15 F #define Z16 D #define Z17 6 #define Z18 1 #define Z19 C #define Z1A 0 #define Z1B 2 #define Z1C B #define Z1D 7 #define Z1E 5 #define Z1F 3 #define Z20 B #define Z21 8 #define Z22 C #define Z23 0 #define Z24 5 #define Z25 2 #define Z26 F #define Z27 D #define Z28 A #define Z29 E #define Z2A 3 #define Z2B 6 #define Z2C 7 #define Z2D 1 #define Z2E 9 #define Z2F 4 #define Z30 7 #define Z31 9 #define Z32 3 #define Z33 1 #define Z34 D #define Z35 C #define Z36 B #define Z37 E #define Z38 2 #define Z39 6 #define Z3A 5 #define Z3B A #define Z3C 4 #define Z3D 0 #define Z3E F #define Z3F 8 #define Z40 9 #define Z41 0 #define Z42 5 #define Z43 7 #define Z44 2 #define Z45 4 #define Z46 A #define Z47 F #define Z48 E #define Z49 1 #define Z4A B #define Z4B C #define Z4C 6 #define Z4D 8 #define Z4E 3 #define Z4F D #define Z50 2 #define Z51 C #define Z52 6 #define Z53 A #define Z54 0 #define Z55 B #define Z56 8 #define Z57 3 #define Z58 4 #define Z59 D #define Z5A 7 #define Z5B 5 #define Z5C F #define Z5D E #define Z5E 1 #define Z5F 9 #define Z60 C #define Z61 5 #define Z62 1 #define Z63 F #define Z64 E #define Z65 D #define Z66 4 #define Z67 A #define Z68 0 #define Z69 7 #define Z6A 6 #define Z6B 3 #define Z6C 9 #define Z6D 2 #define Z6E 8 #define Z6F B #define Z70 D #define Z71 B #define Z72 7 #define Z73 E #define Z74 C #define Z75 1 #define Z76 3 #define Z77 9 #define Z78 5 #define Z79 0 #define Z7A F #define Z7B 4 #define Z7C 8 #define Z7D 6 #define Z7E 2 #define Z7F A #define Z80 6 #define Z81 F #define Z82 E #define Z83 9 #define Z84 B #define Z85 3 #define Z86 0 #define Z87 8 #define Z88 C #define Z89 2 #define Z8A D #define Z8B 7 #define Z8C 1 #define Z8D 4 #define Z8E A #define Z8F 5 #define Z90 A #define Z91 2 #define Z92 8 #define Z93 4 #define Z94 7 #define Z95 6 #define Z96 1 #define Z97 5 #define Z98 F #define Z99 B #define Z9A 9 #define Z9B E #define Z9C 3 #define Z9D C #define Z9E D #define Z9F 0 #define Mx(r, i) Mx_(Z ## r ## i) #define Mx_(n) Mx__(n) #define Mx__(n) M ## n #define CSx(r, i) CSx_(Z ## r ## i) #define CSx_(n) CSx__(n) #define CSx__(n) CS ## n #define CS0 SPH_C32(0x243F6A88) #define CS1 SPH_C32(0x85A308D3) #define CS2 SPH_C32(0x13198A2E) #define CS3 SPH_C32(0x03707344) #define CS4 SPH_C32(0xA4093822) #define CS5 SPH_C32(0x299F31D0) #define CS6 SPH_C32(0x082EFA98) #define CS7 SPH_C32(0xEC4E6C89) #define CS8 SPH_C32(0x452821E6) #define CS9 SPH_C32(0x38D01377) #define CSA SPH_C32(0xBE5466CF) #define CSB SPH_C32(0x34E90C6C) #define CSC SPH_C32(0xC0AC29B7) #define CSD SPH_C32(0xC97C50DD) #define CSE SPH_C32(0x3F84D5B5) #define CSF SPH_C32(0xB5470917) #define GS(m0, m1, c0, c1, a, b, c, d) do { \ a = SPH_T32(a + b + (m0 ^ c1)); \ d = SPH_ROTR32(d ^ a, 16); \ c = SPH_T32(c + d); \ b = SPH_ROTR32(b ^ c, 12); \ a = SPH_T32(a + b + (m1 ^ c0)); \ d = SPH_ROTR32(d ^ a, 8); \ c = SPH_T32(c + d); \ b = SPH_ROTR32(b ^ c, 7); \ } while (0) #define ROUND_S(r) do { \ GS(Mx(r, 0), Mx(r, 1), CSx(r, 0), CSx(r, 1), V0, V4, V8, VC); \ GS(Mx(r, 2), Mx(r, 3), CSx(r, 2), CSx(r, 3), V1, V5, V9, VD); \ GS(Mx(r, 4), Mx(r, 5), CSx(r, 4), CSx(r, 5), V2, V6, VA, VE); \ GS(Mx(r, 6), Mx(r, 7), CSx(r, 6), CSx(r, 7), V3, V7, VB, VF); \ GS(Mx(r, 8), Mx(r, 9), CSx(r, 8), CSx(r, 9), V0, V5, VA, VF); \ GS(Mx(r, A), Mx(r, B), CSx(r, A), CSx(r, B), V1, V6, VB, VC); \ GS(Mx(r, C), Mx(r, D), CSx(r, C), CSx(r, D), V2, V7, V8, VD); \ GS(Mx(r, E), Mx(r, F), CSx(r, E), CSx(r, F), V3, V4, V9, VE); \ } while (0) #define COMPRESS32 do { \ sph_u32 M0, M1, M2, M3, M4, M5, M6, M7; \ sph_u32 M8, M9, MA, MB, MC, MD, ME, MF; \ sph_u32 V0, V1, V2, V3, V4, V5, V6, V7; \ sph_u32 V8, V9, VA, VB, VC, VD, VE, VF; \ V0 = H0; \ V1 = H1; \ V2 = H2; \ V3 = H3; \ V4 = H4; \ V5 = H5; \ V6 = H6; \ V7 = H7; \ V8 = S0 ^ CS0; \ V9 = S1 ^ CS1; \ VA = S2 ^ CS2; \ VB = S3 ^ CS3; \ VC = T0 ^ CS4; \ VD = T0 ^ CS5; \ VE = T1 ^ CS6; \ VF = T1 ^ CS7; \ M0 = input[0]; \ M1 = input[1]; \ M2 = input[2]; \ M3 = input[3]; \ M4 = input[4]; \ M5 = input[5]; \ M6 = input[6]; \ M7 = input[7]; \ M8 = input[8]; \ M9 = input[9]; \ MA = input[10]; \ MB = input[11]; \ MC = input[12]; \ MD = input[13]; \ ME = input[14]; \ MF = input[15]; \ ROUND_S(0); \ ROUND_S(1); \ ROUND_S(2); \ ROUND_S(3); \ ROUND_S(4); \ ROUND_S(5); \ ROUND_S(6); \ ROUND_S(7); \ H0 ^= S0 ^ V0 ^ V8; \ H1 ^= S1 ^ V1 ^ V9; \ H2 ^= S2 ^ V2 ^ VA; \ H3 ^= S3 ^ V3 ^ VB; \ H4 ^= S0 ^ V4 ^ VC; \ H5 ^= S1 ^ V5 ^ VD; \ H6 ^= S2 ^ V6 ^ VE; \ H7 ^= S3 ^ V7 ^ VF; \ } while (0) __global__ void cuda_blake256_hash( uint64_t *g_out, uint32_t nonce, uint32_t *g_good, bool validate ) { uint32_t input[16]; uint64_t output[4]; #pragma unroll for (int i=0; i < 16; ++i) input[i] = pdata[i]; sph_u32 H0 = 0x6A09E667; sph_u32 H1 = 0xBB67AE85; sph_u32 H2 = 0x3C6EF372; sph_u32 H3 = 0xA54FF53A; sph_u32 H4 = 0x510E527F; sph_u32 H5 = 0x9B05688C; sph_u32 H6 = 0x1F83D9AB; sph_u32 H7 = 0x5BE0CD19; sph_u32 S0 = 0; sph_u32 S1 = 0; sph_u32 S2 = 0; sph_u32 S3 = 0; sph_u32 T0 = 0; sph_u32 T1 = 0; T0 = SPH_T32(T0 + 512); COMPRESS32; #pragma unroll for (int i=0; i < 3; ++i) input[i] = pdata[16+i]; input[3] = nonce + ((blockIdx.x * blockDim.x) + threadIdx.x); input[4] = 0x80000000; #pragma unroll 8 for (int i=5; i < 13; ++i) input[i] = 0; input[13] = 0x00000001; input[14] = T1; input[15] = T0 + 128; T0 = SPH_T32(T0 + 128); COMPRESS32; cuda_sph_enc32be((unsigned char*)output + 4*6, H6); cuda_sph_enc32be((unsigned char*)output + 4*7, H7); if (validate || output[3] <= ptarget64[3]) { // this data is only needed when we actually need to save the hashes cuda_sph_enc32be((unsigned char*)output + 4*0, H0); cuda_sph_enc32be((unsigned char*)output + 4*1, H1); cuda_sph_enc32be((unsigned char*)output + 4*2, H2); cuda_sph_enc32be((unsigned char*)output + 4*3, H3); cuda_sph_enc32be((unsigned char*)output + 4*4, H4); cuda_sph_enc32be((unsigned char*)output + 4*5, H5); } if (validate) { g_out += 4 * ((blockIdx.x * blockDim.x) + threadIdx.x); #pragma unroll for (int i=0; i < 4; ++i) g_out[i] = output[i]; } if (output[3] <= ptarget64[3]) { uint64_t *g_good64 = (uint64_t*)g_good; if (output[3] < g_good64[3]) { g_good64[3] = output[3]; g_good64[2] = output[2]; g_good64[1] = output[1]; g_good64[0] = output[0]; g_good[8] = nonce + ((blockIdx.x * blockDim.x) + threadIdx.x); } } } static std::map<int, uint32_t *> context_good[2]; static bool init[MAX_GPUS] = { 0 }; bool default_prepare_blake256(int thr_id, const uint32_t host_pdata[20], const uint32_t host_ptarget[8]) { if (!init[thr_id]) { // allocate pinned host memory for good hashes uint32_t *tmp; checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[0][thr_id] = tmp; checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[1][thr_id] = tmp; init[thr_id] = true; } checkCudaErrors(cudaMemcpyToSymbol(pdata, host_pdata, 80, 0, cudaMemcpyHostToDevice)); checkCudaErrors(cudaMemcpyToSymbol(ptarget64, host_ptarget, 32, 0, cudaMemcpyHostToDevice)); return context_good[0][thr_id] && context_good[1][thr_id]; } void default_do_blake256(dim3 grid, dim3 threads, int thr_id, int stream, uint32_t *hash, uint32_t nonce, int throughput, bool do_d2h) { checkCudaErrors(cudaMemsetAsync(context_good[stream][thr_id], 0xff, 9 * sizeof(uint32_t), context_streams[stream][thr_id])); cuda_blake256_hash<<<grid, threads, 0, context_streams[stream][thr_id]>>>((uint64_t*)context_hash[stream][thr_id], nonce, context_good[stream][thr_id], do_d2h); // copy hashes from device memory to host (ALL hashes, lots of data...) if (do_d2h && hash != NULL) { size_t mem_size = throughput * sizeof(uint32_t) * 8; checkCudaErrors(cudaMemcpyAsync(hash, context_hash[stream][thr_id], mem_size, cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); } else if (hash != NULL) { // asynchronous copy of winning nonce (just 4 bytes...) checkCudaErrors(cudaMemcpyAsync(hash, context_good[stream][thr_id]+8, sizeof(uint32_t), cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); } } void default_free_blake256(int thr_id) { if (init[thr_id]) { cudaFree(context_good[0][thr_id]); cudaFree(context_good[1][thr_id]); init[thr_id] = false; } }