/* * sha256 djm34, catia * */ /* * sha-256 kernel implementation. * * ==========================(LICENSE BEGIN)============================ * * Copyright (c) 2014 djm34 * * Permission is hereby granted, free of charge, to any person obtaining * a copy of this software and associated documentation files (the * "Software"), to deal in the Software without restriction, including * without limitation the rights to use, copy, modify, merge, publish, * distribute, sublicense, and/or sell copies of the Software, and to * permit persons to whom the Software is furnished to do so, subject to * the following conditions: * * The above copyright notice and this permission notice shall be * included in all copies or substantial portions of the Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. * * ===========================(LICENSE END)============================= * * @author phm */ #undef _GLIBCXX_ATOMIC_BUILTINS #undef _GLIBCXX_USE_INT128 #include #include //#include "uint256.h" #include "cuda_helper.h" extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); //#include "cuPrintf.cu" typedef struct t4_t{ uint64_t high,low; } t4_t; __device__ __forceinline__ ulonglong2 umul64wide(unsigned long long int a, unsigned long long int b) { ulonglong2 res; asm ("{\n\t" ".reg .u32 r0, r1, r2, r3, alo, ahi, blo, bhi;\n\t" "mov.b64 {alo,ahi}, %2; \n\t" "mov.b64 {blo,bhi}, %3; \n\t" "mul.lo.u32 r0, alo, blo; \n\t" "mul.hi.u32 r1, alo, blo; \n\t" "mad.lo.cc.u32 r1, alo, bhi, r1;\n\t" "madc.hi.u32 r2, alo, bhi, 0;\n\t" "mad.lo.cc.u32 r1, ahi, blo, r1;\n\t" "madc.hi.cc.u32 r2, ahi, blo, r2;\n\t" "madc.hi.u32 r3, ahi, bhi, 0;\n\t" "mad.lo.cc.u32 r2, ahi, bhi, r2;\n\t" "addc.u32 r3, r3, 0; \n\t" "mov.b64 %0, {r0,r1}; \n\t" "mov.b64 %1, {r2,r3}; \n\t" "}" : "=l"(res.x), "=l"(res.y) : "l"(a), "l"(b)); return res; } #define umul_ppmm(h,l,m,n) \ { \ ulonglong2 foom = umul64wide(m,n); \ h = foom.y; \ l = foom.x; \ } __device__ __forceinline__ void umul_ppmmT4(t4_t *h, t4_t *l, t4_t m, t4_t n) { asm ("{\n\t" ".reg .u32 o0, o1, o2, o3, o4; \n\t" ".reg .u32 o5, o6, o7, i8, i9; \n\t" ".reg .u32 i10, i11, i12, i13; \n\t" ".reg .u32 i14, i15, i16, i17; \n\t" ".reg .u32 i18, i19, i20, i21; \n\t" ".reg .u32 i22, i23; \n\t" "mov.b64 { i8, i9}, %4; \n\t" "mov.b64 {i10,i11}, %5; \n\t" "mov.b64 {i12,i13}, %6; \n\t" "mov.b64 {i14,i15}, %7; \n\t" "mov.b64 {i16,i17}, %8; \n\t" "mov.b64 {i18,i19}, %9; \n\t" "mov.b64 {i20,i21},%10; \n\t" "mov.b64 {i22,i23},%11; \n\t" "mul.lo.u32 o0, i8, i16; \n\t" "mul.hi.u32 o1, i8, i16; \n\t" "mad.lo.cc.u32 o1, i8, i17, o1;\n\t" "madc.hi.u32 o2, i8, i17, 0;\n\t" "mad.lo.cc.u32 o1, i9, i16, o1;\n\t" "madc.hi.cc.u32 o2, i9, i16, o2;\n\t" "madc.hi.u32 o3, i8, i18, 0;\n\t" "mad.lo.cc.u32 o2, i8, i18, o2;\n\t" "madc.hi.cc.u32 o3, i9, i17, o3;\n\t" "madc.hi.u32 o4, i8, i19, 0;\n\t" "mad.lo.cc.u32 o2, i9, i17, o2;\n\t" "madc.hi.cc.u32 o3, i10, i16, o3;\n\t" "madc.hi.cc.u32 o4, i9, i18, o4;\n\t" "addc.u32 o5, 0, 0;\n\t" "mad.lo.cc.u32 o2, i10, i16, o2;\n\t" "madc.lo.cc.u32 o3, i8, i19, o3;\n\t" "madc.hi.cc.u32 o4, i10, i17, o4;\n\t" "madc.hi.cc.u32 o5, i9, i19, o5;\n\t" "addc.u32 o6, 0, 0;\n\t" "mad.lo.cc.u32 o3, i9, i18, o3;\n\t" "madc.hi.cc.u32 o4, i11, i16, o4;\n\t" "madc.hi.cc.u32 o5, i10, i18, o5;\n\t" "addc.u32 o6, 0, o6;\n\t" "mad.lo.cc.u32 o3, i10, i17, o3;\n\t" "addc.u32 o4, 0, o4;\n\t" "mad.hi.cc.u32 o5, i11, i17, o5;\n\t" "madc.hi.cc.u32 o6, i10, i19, o6;\n\t" "addc.u32 o7, 0, 0;\n\t" "mad.lo.cc.u32 o3, i11, i16, o3;\n\t" "madc.lo.cc.u32 o4, i9, i19, o4;\n\t" "addc.u32 o5, 0, o5;\n\t" "mad.hi.cc.u32 o6, i11, i18, o6;\n\t" "addc.u32 o7, 0, o7;\n\t" "mad.lo.cc.u32 o4, i10, i18, o4;\n\t" "addc.u32 o5, 0, o5;\n\t" "mad.hi.u32 o7, i11, i19, o7;\n\t" "mad.lo.cc.u32 o4, i11, i17, o4;\n\t" "addc.u32 o5, 0, o5;\n\t" "mad.lo.cc.u32 o5, i10, i19, o5;\n\t" "addc.u32 o6, 0, o6;\n\t" "mad.lo.cc.u32 o5, i11, i18, o5;\n\t" "addc.u32 o6, 0, o6;\n\t" "mad.lo.cc.u32 o6, i11, i19, o6;\n\t" "addc.u32 o7, 0, o7;\n\t" "mov.b64 %0, {o0,o1}; \n\t" "mov.b64 %1, {o2,o3}; \n\t" "mov.b64 %2, {o4,o5}; \n\t" "mov.b64 %3, {o6,o7}; \n\t" "}" : "=l"(l->low), "=l"(l->high), "=l"(h->low), "=l"(h->high) : "l"(m.low), "l"(m.high), "l"(0ULL), "l"(0ULL), "l"(n.low), "l"(n.high), "l"(0ULL), "l"(0ULL)); } #if 0 __device__ __forceinline__ void umul_ppmmT4(t4_t *h, t4_t *l, t4_t m, t4_t n){ uint64_t th,tl; uint32_t c,c2; umul_ppmm(l->high,l->low,m.low,n.low); umul_ppmm(th,tl,m.high,n.low); l->high += tl; c = (l->high < tl); h->low = th + c; c = (h->low < c); h->high = c; //Second word umul_ppmm(th,tl,m.low,n.high); l->high += tl; c = l->high < tl; h->low += th; c2 = h->low < th; h->low += c; c2 += h->low < c; h->high += c2; umul_ppmm(th,tl,m.high,n.high); h->low += tl; c = h->low < tl; h->high += th + c; } #endif __device__ __forceinline__ t4_t T4(uint32_t thread, uint32_t threads, uint32_t idx, uint64_t *g){ t4_t ret; ret.high = g[(idx*2 + 1)*threads + thread]; ret.low = g[(idx*2)*threads + thread]; if(thread==0){ // cuPrintf("Load Idx: %d %8.8X %8.8X %8.8X %8.8X\n", idx, ret.high>>32, ret.high, ret.low>>32, ret.low); } return ret; } __device__ __forceinline__ void T4_store(uint32_t thread, uint32_t threads, uint32_t idx, uint64_t *g, t4_t val){ g[(idx*2 + 1)*threads + thread]=val.high; g[(idx*2)*threads + thread]=val.low; if(thread==0){ // cuPrintf("Store Idx: %d %8.8X %8.8X %8.8X %8.8X\n", idx, val.high>>32, val.high, val.low>>32, val.low); } } __device__ __forceinline__ void T4_set(t4_t *d, uint64_t v){ d->high = 0; d->low = v; } __device__ __forceinline__ t4_t T4_add(t4_t a, t4_t b){ t4_t ret; uint32_t c=0; ret.low = a.low + b.low; if(ret.low < a.low) c=1; ret.high = a.high + b.high + c; return ret; } __device__ __forceinline__ t4_t T4_add(uint64_t a, t4_t b){ t4_t ret; uint32_t c=0; ret.low = a + b.low; if(ret.low < a) c=1; ret.high = b.high + c; return ret; } __device__ __forceinline__ uint32_t T4_lt(t4_t a, t4_t b){ if(a.high < b.high) return 1; if(a.high == b.high && a.low < b.low) return 1; return 0; } __device__ __forceinline__ uint32_t T4_gt(t4_t a, uint64_t b){ if(a.high) return 1; if(a.low > b) return 1; return 0; } __device__ void mulScalarT4(uint32_t thread, uint32_t threads, uint32_t len, uint64_t* g_p, uint64_t* g_v, t4_t sml, uint32_t *size){ t4_t ul, cl, hpl, lpl; uint32_t i; T4_set(&cl,0); for(i=0; i < len; i++) { ul = T4(thread,threads,i,g_v); umul_ppmmT4 (&hpl, &lpl, ul, sml); lpl = T4_add(lpl,cl); cl = T4_add(T4_lt(lpl,cl),hpl); T4_store(thread,threads,i,g_p,lpl); } T4_store(thread,threads,len,g_p,cl); *size = len + T4_gt(cl,0); } __device__ void mulScalar(uint32_t thread, uint32_t threads, uint32_t len, uint64_t* g_p, uint64_t* g_v, uint64_t sml, uint32_t *size){ uint64_t ul, cl, hpl, lpl; uint32_t i; cl = 0; for(i=0; i < len; i++) { ul = g_v[i*threads + thread]; umul_ppmm (hpl, lpl, ul, sml); lpl += cl; cl = (lpl < cl) + hpl; g_p[i*threads + thread] = lpl; } g_p[len*threads + thread] = cl; *size = len + (cl != 0); } uint64_t __device__ addmul_1g (uint32_t thread, uint32_t threads, uint64_t *sum, uint32_t sofst, uint64_t *x, uint64_t xsz, uint64_t a) { uint64_t carry=0; uint32_t i; uint64_t ul,lpl,hpl,rl; for(i=0; i < xsz; i++) { ul = x[i*threads + thread]; umul_ppmm (hpl, lpl, ul, a); lpl += carry; carry = (lpl < carry) + hpl; rl = sum[(i+sofst) * threads + thread]; lpl = rl + lpl; carry += lpl < rl; sum[(i+sofst)*threads + thread] = lpl; } return carry; } __device__ t4_t addmul_1gT4 (uint32_t thread, uint32_t threads, uint64_t *sum, uint32_t sofst, uint64_t *x, uint64_t xsz, t4_t a) { t4_t carry; uint32_t i; t4_t ul,lpl,hpl,rl; T4_set(&carry,0); for(i=0; i < xsz; i++) { ul = T4(thread,threads,i,x); umul_ppmmT4 (&hpl, &lpl, ul, a); lpl = T4_add(lpl,carry); carry = T4_add(T4_lt(lpl,carry), hpl); rl = T4(thread,threads,i+sofst,sum); lpl = T4_add(rl,lpl); carry = T4_add(T4_lt(lpl,rl),carry); T4_store(thread,threads,i+sofst,sum,lpl); } return carry; } __global__ void gpu_mul(int threads, uint32_t ulegs, uint32_t vlegs, uint64_t *g_u, uint64_t *g_v, uint64_t *g_p) { int thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) { if(ulegs < vlegs) { uint64_t t1=ulegs; ulegs = vlegs; vlegs = t1; uint64_t *t2 = g_u; g_u = g_v; g_v = t2; } uint32_t vofst=1,rofst=1,psize=0; mulScalar(thread,threads,ulegs,g_p,g_u,g_v[thread],&psize); #if 1 while (vofst < vlegs) { //clear high word //TODO: right // printf("Size: %d\n", rp->size[tid]); g_p[(psize+0)*threads + thread] = 0; g_p[(ulegs+rofst)*threads + thread] = addmul_1g (thread, threads, g_p ,rofst , g_u, ulegs, g_v[vofst*threads+thread]); vofst++; rofst++; psize++; } // if(D_REF(rp->d,up->size[tid] + vp->size[tid] - 1,tid) != (uint64_t)0) // rp->size[tid]++; #endif } } __global__ void gpu_mulT4(int threads, uint32_t ulegs, uint32_t vlegs, uint64_t *g_u, uint64_t *g_v, uint64_t *g_p) { int thread = (blockDim.x * blockIdx.x + threadIdx.x); if (thread < threads) { if(ulegs < vlegs) { // everything written the other way around... are you kidding me ?! uint64_t t1=ulegs; ulegs = vlegs; vlegs = t1; uint64_t *t2 = g_u; g_u = g_v; g_v = t2; } ulegs >>= 1; vlegs >>= 1; if(thread == 0) { // cuPrintf("U: %d V: %d\n", ulegs, vlegs); } uint32_t vofst=1,rofst=1,psize=0; mulScalarT4(thread,threads,ulegs,g_p,g_u,T4(thread,threads,0,g_v),&psize); #if 1 t4_t zero; T4_set(&zero,0); // while (vofst < vlegs) { #pragma unroll for (vofst=1;vofst>>(threads, alegs, blegs, g_a, g_b, g_p) ; } __host__ void cpu_mulT4(int thr_id, int threads, uint32_t alegs, uint32_t blegs, uint64_t *g_a, uint64_t *g_b, uint64_t *g_p, int order) { const int threadsperblock = 256; // better occupancy (for both 780 and 750 ti's) // berechne wie viele Thread Blocks wir brauchen dim3 grid((threads + threadsperblock-1)/threadsperblock); dim3 block(threadsperblock); size_t shared_size = 0; //gpu_mulT4<<>>(threads, alegs, blegs, g_a, g_b, g_p) ; gpu_mulT4<<>>(threads, blegs, alegs, g_b, g_a, g_p) ; } __host__ void mul_init() { }