GOSTSec
/
ccminer-gostd-lite
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
GOSTCoin CUDA miner project, compatible with most nvidia cards, containing only gostd algo
You can not select more than 25 topics Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
766 Commits
1 Branch
55 Tags
65 MiB
 
 
 
 
 
 
Tree: 35bfe8d191
Branches Tags
${ item.name }
Create tag ${ searchTerm }
Create branch ${ searchTerm }
from '35bfe8d191'
${ noResults }
ccminer-gostd-lite/scrypt/fermi_kernel.cu

907 lines
33 KiB
Raw Blame History

//
// Kernel that runs best on Fermi devices
//
// - shared memory use reduced by nearly factor 2 over legacy kernel
// by transferring only half work units (16 x uint32_t) at once.
// - uses ulong2/uint4 based memory transfers (each thread moves 16 bytes),
// allowing for shorter unrolled loops. This relies on Fermi's better
// memory controllers to get high memory troughput.
//
// NOTE: compile this .cu module for compute_20,sm_20 with --maxrregcount=63
//
// TODO: batch-size support for this kernel
//
#include <map>
#include <cuda_runtime.h>
#include "miner.h"
#include "salsa_kernel.h"
#include "fermi_kernel.h"
#define THREADS_PER_WU 1 // single thread per hash
#define TEXWIDTH 32768
// forward references
template <int ALGO> __global__ void fermi_scrypt_core_kernelA(uint32_t *g_idata, unsigned int N);
template <int ALGO> __global__ void fermi_scrypt_core_kernelB(uint32_t *g_odata, unsigned int N);
template <int ALGO, int TEX_DIM> __global__ void fermi_scrypt_core_kernelB_tex(uint32_t *g_odata, unsigned int N);
template <int ALGO> __global__ void fermi_scrypt_core_kernelA_LG(uint32_t *g_idata, unsigned int N, unsigned int LOOKUP_GAP);
template <int ALGO> __global__ void fermi_scrypt_core_kernelB_LG(uint32_t *g_odata, unsigned int N, unsigned int LOOKUP_GAP);
template <int ALGO, int TEX_DIM> __global__ void fermi_scrypt_core_kernelB_LG_tex(uint32_t *g_odata, unsigned int N, unsigned int LOOKUP_GAP);
// scratchbuf constants (pointers to scratch buffer for each warp, i.e. 32 hashes)
__constant__ uint32_t* c_V[TOTAL_WARP_LIMIT];
// using texture references for the "tex" variants of the B kernels
texture<uint4, 1, cudaReadModeElementType> texRef1D_4_V;
texture<uint4, 2, cudaReadModeElementType> texRef2D_4_V;
FermiKernel::FermiKernel() : KernelInterface()
{
}
bool FermiKernel::bindtexture_1D(uint32_t *d_V, size_t size)
{
cudaChannelFormatDesc channelDesc4 = cudaCreateChannelDesc<uint4>();
texRef1D_4_V.normalized = 0;
texRef1D_4_V.filterMode = cudaFilterModePoint;
texRef1D_4_V.addressMode[0] = cudaAddressModeClamp;
checkCudaErrors(cudaBindTexture(NULL, &texRef1D_4_V, d_V, &channelDesc4, size));
return true;
}
bool FermiKernel::bindtexture_2D(uint32_t *d_V, int width, int height, size_t pitch)
{
cudaChannelFormatDesc channelDesc4 = cudaCreateChannelDesc<uint4>();
texRef2D_4_V.normalized = 0;
texRef2D_4_V.filterMode = cudaFilterModePoint;
texRef2D_4_V.addressMode[0] = cudaAddressModeClamp;
texRef2D_4_V.addressMode[1] = cudaAddressModeClamp;
// maintain texture width of TEXWIDTH (max. limit is 65000)
while (width > TEXWIDTH) { width /= 2; height *= 2; pitch /= 2; }
while (width < TEXWIDTH) { width *= 2; height = (height+1)/2; pitch *= 2; }
// fprintf(stderr, "total size: %u, %u bytes\n", pitch * height, width * sizeof(uint32_t) * 4 * height);
// fprintf(stderr, "binding width width=%d, height=%d, pitch=%d\n", width, height,pitch);
checkCudaErrors(cudaBindTexture2D(NULL, &texRef2D_4_V, d_V, &channelDesc4, width, height, pitch));
return true;
}
bool FermiKernel::unbindtexture_1D()
{
checkCudaErrors(cudaUnbindTexture(texRef1D_4_V));
return true;
}
bool FermiKernel::unbindtexture_2D()
{
checkCudaErrors(cudaUnbindTexture(texRef2D_4_V));
return true;
}
void FermiKernel::set_scratchbuf_constants(int MAXWARPS, uint32_t** h_V)
{
checkCudaErrors(cudaMemcpyToSymbol(c_V, h_V, MAXWARPS*sizeof(uint32_t*), 0, cudaMemcpyHostToDevice));
}
bool FermiKernel::run_kernel(dim3 grid, dim3 threads, int WARPS_PER_BLOCK, int thr_id, cudaStream_t stream, uint32_t* d_idata, uint32_t* d_odata, unsigned int N, unsigned int LOOKUP_GAP, bool interactive, bool benchmark, int texture_cache)
{
bool success = true;
int shared = WARPS_PER_BLOCK * WU_PER_WARP * (16+4) * sizeof(uint32_t);
// First phase: Sequential writes to scratchpad.
if (LOOKUP_GAP == 1) {
if (IS_SCRYPT()) fermi_scrypt_core_kernelA<A_SCRYPT><<< grid, threads, shared, stream >>>(d_idata, N);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelA<A_SCRYPT_JANE><<< grid, threads, shared, stream >>>(d_idata, N);
} else {
if (IS_SCRYPT()) fermi_scrypt_core_kernelA_LG<A_SCRYPT><<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelA_LG<A_SCRYPT_JANE><<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP);
}
// Second phase: Random read access from scratchpad.
if (LOOKUP_GAP == 1) {
if (texture_cache) {
if (texture_cache == 1) {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB_tex<A_SCRYPT,1><<< grid, threads, shared, stream >>>(d_odata, N);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB_tex<A_SCRYPT_JANE,1><<< grid, threads, shared, stream >>>(d_odata, N);
} else if (texture_cache == 2) {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB_tex<A_SCRYPT,2><<< grid, threads, shared, stream >>>(d_odata, N);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB_tex<A_SCRYPT_JANE,2><<< grid, threads, shared, stream >>>(d_odata, N);
}
else success = false;
} else {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB<A_SCRYPT><<< grid, threads, shared, stream >>>(d_odata, N);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB<A_SCRYPT_JANE><<< grid, threads, shared, stream >>>(d_odata, N);
}
} else {
if (texture_cache) {
if (texture_cache == 1) {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB_LG_tex<A_SCRYPT,1><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB_LG_tex<A_SCRYPT_JANE,1><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
} else if (texture_cache == 2) {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB_LG_tex<A_SCRYPT,2><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB_LG_tex<A_SCRYPT_JANE,2><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
}
else success = false;
} else {
if (IS_SCRYPT()) fermi_scrypt_core_kernelB_LG<A_SCRYPT><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
if (IS_SCRYPT_JANE()) fermi_scrypt_core_kernelB_LG<A_SCRYPT_JANE><<< grid, threads, shared, stream >>>(d_odata, N, LOOKUP_GAP);
}
}
return success;
}
#if 0
#define ROTL(a, b) (((a) << (b)) | ((a) >> (32 - (b))))
#define QUARTER(a,b,c,d) \
a += b; d ^= a; d = ROTL(d,16); \
c += d; b ^= c; b = ROTL(b,12); \
a += b; d ^= a; d = ROTL(d,8); \
c += d; b ^= c; b = ROTL(b,7);
static __device__ void xor_chacha8(uint4 *B, uint4 *C)
{
uint32_t x[16];
x[0]=(B[0].x ^= C[0].x);
x[1]=(B[0].y ^= C[0].y);
x[2]=(B[0].z ^= C[0].z);
x[3]=(B[0].w ^= C[0].w);
x[4]=(B[1].x ^= C[1].x);
x[5]=(B[1].y ^= C[1].y);
x[6]=(B[1].z ^= C[1].z);
x[7]=(B[1].w ^= C[1].w);
x[8]=(B[2].x ^= C[2].x);
x[9]=(B[2].y ^= C[2].y);
x[10]=(B[2].z ^= C[2].z);
x[11]=(B[2].w ^= C[2].w);
x[12]=(B[3].x ^= C[3].x);
x[13]=(B[3].y ^= C[3].y);
x[14]=(B[3].z ^= C[3].z);
x[15]=(B[3].w ^= C[3].w);
/* Operate on columns. */
QUARTER( x[0], x[4], x[ 8], x[12] )
QUARTER( x[1], x[5], x[ 9], x[13] )
QUARTER( x[2], x[6], x[10], x[14] )
QUARTER( x[3], x[7], x[11], x[15] )
/* Operate on diagonals */
QUARTER( x[0], x[5], x[10], x[15] )
QUARTER( x[1], x[6], x[11], x[12] )
QUARTER( x[2], x[7], x[ 8], x[13] )
QUARTER( x[3], x[4], x[ 9], x[14] )
/* Operate on columns. */
QUARTER( x[0], x[4], x[ 8], x[12] )
QUARTER( x[1], x[5], x[ 9], x[13] )
QUARTER( x[2], x[6], x[10], x[14] )
QUARTER( x[3], x[7], x[11], x[15] )
/* Operate on diagonals */
QUARTER( x[0], x[5], x[10], x[15] )
QUARTER( x[1], x[6], x[11], x[12] )
QUARTER( x[2], x[7], x[ 8], x[13] )
QUARTER( x[3], x[4], x[ 9], x[14] )
/* Operate on columns. */
QUARTER( x[0], x[4], x[ 8], x[12] )
QUARTER( x[1], x[5], x[ 9], x[13] )
QUARTER( x[2], x[6], x[10], x[14] )
QUARTER( x[3], x[7], x[11], x[15] )
/* Operate on diagonals */
QUARTER( x[0], x[5], x[10], x[15] )
QUARTER( x[1], x[6], x[11], x[12] )
QUARTER( x[2], x[7], x[ 8], x[13] )
QUARTER( x[3], x[4], x[ 9], x[14] )
/* Operate on columns. */
QUARTER( x[0], x[4], x[ 8], x[12] )
QUARTER( x[1], x[5], x[ 9], x[13] )
QUARTER( x[2], x[6], x[10], x[14] )
QUARTER( x[3], x[7], x[11], x[15] )
/* Operate on diagonals */
QUARTER( x[0], x[5], x[10], x[15] )
QUARTER( x[1], x[6], x[11], x[12] )
QUARTER( x[2], x[7], x[ 8], x[13] )
QUARTER( x[3], x[4], x[ 9], x[14] )
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7];
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15];
}
#else
#define ROTL(a, b) (((a) << (b)) | ((a) >> (32 - (b))))
#define ADD4(d1,d2,d3,d4,s1,s2,s3,s4) \
d1 += s1; d2 += s2; d3 += s3; d4 += s4;
#define XOR4(d1,d2,d3,d4,s1,s2,s3,s4) \
d1 ^= s1; d2 ^= s2; d3 ^= s3; d4 ^= s4;
#define ROTL4(d1,d2,d3,d4,amt) \
d1 = ROTL(d1, amt); d2 = ROTL(d2, amt); d3 = ROTL(d3, amt); d4 = ROTL(d4, amt);
#define QROUND(a1,a2,a3,a4, b1,b2,b3,b4, c1,c2,c3,c4, amt) \
ADD4 (a1,a2,a3,a4, c1,c2,c3,c4) \
XOR4 (b1,b2,b3,b4, a1,a2,a3,a4) \
ROTL4(b1,b2,b3,b4, amt)
static __device__ void xor_chacha8(uint4 *B, uint4 *C)
{
uint32_t x[16];
x[0]=(B[0].x ^= C[0].x);
x[1]=(B[0].y ^= C[0].y);
x[2]=(B[0].z ^= C[0].z);
x[3]=(B[0].w ^= C[0].w);
x[4]=(B[1].x ^= C[1].x);
x[5]=(B[1].y ^= C[1].y);
x[6]=(B[1].z ^= C[1].z);
x[7]=(B[1].w ^= C[1].w);
x[8]=(B[2].x ^= C[2].x);
x[9]=(B[2].y ^= C[2].y);
x[10]=(B[2].z ^= C[2].z);
x[11]=(B[2].w ^= C[2].w);
x[12]=(B[3].x ^= C[3].x);
x[13]=(B[3].y ^= C[3].y);
x[14]=(B[3].z ^= C[3].z);
x[15]=(B[3].w ^= C[3].w);
/* Operate on columns. */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7);
/* Operate on diagonals */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7);
/* Operate on columns. */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7);
/* Operate on diagonals */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7);
/* Operate on columns. */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7);
/* Operate on diagonals */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7);
/* Operate on columns. */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8);
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7);
/* Operate on diagonals */
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12);
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8);
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7);
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7];
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15];
}
#endif
#define ROTL7(a0,a1,a2,a3,a00,a10,a20,a30){\
a0^=(((a00)<<7) | ((a00)>>25) );\
a1^=(((a10)<<7) | ((a10)>>25) );\
a2^=(((a20)<<7) | ((a20)>>25) );\
a3^=(((a30)<<7) | ((a30)>>25) );\
};\
#define ROTL9(a0,a1,a2,a3,a00,a10,a20,a30){\
a0^=(((a00)<<9) | ((a00)>>23) );\
a1^=(((a10)<<9) | ((a10)>>23) );\
a2^=(((a20)<<9) | ((a20)>>23) );\
a3^=(((a30)<<9) | ((a30)>>23) );\
};\
#define ROTL13(a0,a1,a2,a3,a00,a10,a20,a30){\
a0^=(((a00)<<13) | ((a00)>>19) );\
a1^=(((a10)<<13) | ((a10)>>19) );\
a2^=(((a20)<<13) | ((a20)>>19) );\
a3^=(((a30)<<13) | ((a30)>>19) );\
};\
#define ROTL18(a0,a1,a2,a3,a00,a10,a20,a30){\
a0^=(((a00)<<18) | ((a00)>>14) );\
a1^=(((a10)<<18) | ((a10)>>14) );\
a2^=(((a20)<<18) | ((a20)>>14) );\
a3^=(((a30)<<18) | ((a30)>>14) );\
};\
static __device__ void xor_salsa8(uint4 *B, uint4 *C)
{
uint32_t x[16];
x[0]=(B[0].x ^= C[0].x);
x[1]=(B[0].y ^= C[0].y);
x[2]=(B[0].z ^= C[0].z);
x[3]=(B[0].w ^= C[0].w);
x[4]=(B[1].x ^= C[1].x);
x[5]=(B[1].y ^= C[1].y);
x[6]=(B[1].z ^= C[1].z);
x[7]=(B[1].w ^= C[1].w);
x[8]=(B[2].x ^= C[2].x);
x[9]=(B[2].y ^= C[2].y);
x[10]=(B[2].z ^= C[2].z);
x[11]=(B[2].w ^= C[2].w);
x[12]=(B[3].x ^= C[3].x);
x[13]=(B[3].y ^= C[3].y);
x[14]=(B[3].z ^= C[3].z);
x[15]=(B[3].w ^= C[3].w);
/* Operate on columns. */
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]);
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]);
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]);
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]);
/* Operate on rows. */
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]);
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]);
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]);
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]);
/* Operate on columns. */
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]);
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]);
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]);
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]);
/* Operate on rows. */
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]);
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]);
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]);
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]);
/* Operate on columns. */
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]);
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]);
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]);
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]);
/* Operate on rows. */
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]);
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]);
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]);
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]);
/* Operate on columns. */
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]);
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]);
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]);
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]);
/* Operate on rows. */
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]);
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]);
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]);
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]);
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7];
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15];
}
static __device__ __forceinline__ uint4& operator^=(uint4& left, const uint4& right)
{
left.x ^= right.x;
left.y ^= right.y;
left.z ^= right.z;
left.w ^= right.w;
return left;
}
////////////////////////////////////////////////////////////////////////////////
//! Scrypt core kernel for Fermi class devices.
//! @param g_idata input data in global memory
//! @param g_odata output data in global memory
////////////////////////////////////////////////////////////////////////////////
template <int ALGO> __global__
void fermi_scrypt_core_kernelA(uint32_t *g_idata, unsigned int N)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
const unsigned int LOOKUP_GAP = 1;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_idata += 32 * offset;
uint32_t * V = c_V[offset / WU_PER_WARP] + SCRATCH*Y + Z;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu])) = *((ulonglong2*)XB[wu]) = *((ulonglong2*)(&g_idata[32*(wu+Y)+Z]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu+16])) = *((ulonglong2*)XB[wu]) = *((ulonglong2*)(&g_idata[32*(wu+Y)+16+Z]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
for (int i = 1; i < N; i++) {
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu + i*32])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu + i*32 + 16])) = *((ulonglong2*)XB[wu]);
}
}
template <int ALGO> __global__
void fermi_scrypt_core_kernelB(uint32_t *g_odata, unsigned int N)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
const unsigned int LOOKUP_GAP = 1;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_odata += 32 * offset;
uint32_t * V = c_V[offset / WU_PER_WARP] + SCRATCH*Y + Z;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + (N-1)*32]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + (N-1)*32 + 16]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
for (int i = 0; i < N; i++) {
XX[16] = 32 * (C[0].x & (N-1));
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + XB[wu][16-Z]]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] ^= *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + XB[wu][16-Z] + 16]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] ^= *((uint4*)&XX[4*idx]);
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+Z])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+16+Z])) = *((ulonglong2*)XB[wu]);
}
template <int ALGO, int TEX_DIM> __global__ void
fermi_scrypt_core_kernelB_tex(uint32_t *g_odata, unsigned int N)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
const unsigned int LOOKUP_GAP = 1;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_odata += 32 * offset;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + (N-1)*32 + Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + (N-1)*32 + 16+Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
for (int i = 0; i < N; i++) {
XX[16] = 32 * (C[0].x & (N-1));
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + XB[wu][16-Z] + Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] ^= *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + XB[wu][16-Z] + 16+Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] ^= *((uint4*)&XX[4*idx]);
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+Z])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+16+Z])) = *((ulonglong2*)XB[wu]);
}
//
// Lookup-Gap variations of the above functions
//
template <int ALGO> __global__ void
fermi_scrypt_core_kernelA_LG(uint32_t *g_idata, unsigned int N, unsigned int LOOKUP_GAP)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_idata += 32 * offset;
uint32_t * V = c_V[offset / WU_PER_WARP] + SCRATCH*Y + Z;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu])) = *((ulonglong2*)XB[wu]) = *((ulonglong2*)(&g_idata[32*(wu+Y)+Z]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu+16])) = *((ulonglong2*)XB[wu]) = *((ulonglong2*)(&g_idata[32*(wu+Y)+16+Z]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
for (int i = 1; i < N; i++) {
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
if (i % LOOKUP_GAP == 0) {
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu + (i/LOOKUP_GAP)*32])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&V[SCRATCH*wu + (i/LOOKUP_GAP)*32 + 16])) = *((ulonglong2*)XB[wu]);
}
}
}
template <int ALGO> __global__ void
fermi_scrypt_core_kernelB_LG(uint32_t *g_odata, unsigned int N, unsigned int LOOKUP_GAP)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_odata += 32 * offset;
uint32_t * V = c_V[offset / WU_PER_WARP] + SCRATCH*Y + Z;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
uint32_t pos = (N-1)/LOOKUP_GAP; uint32_t loop = 1 + (N-1)-pos*LOOKUP_GAP;
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + pos*32]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + pos*32 + 16]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
while (loop--)
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
for (int i = 0; i < N; i++) {
uint32_t j = C[0].x & (N-1);
uint32_t pos = j / LOOKUP_GAP; uint32_t loop = j - pos*LOOKUP_GAP;
XX[16] = 32 * pos;
uint4 b[4], c[4];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + XB[wu][16-Z]]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) b[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)XB[wu]) = *((ulonglong2*)(&V[SCRATCH*wu + XB[wu][16-Z] + 16]));
#pragma unroll 4
for (int idx=0; idx < 4; idx++) c[idx] = *((uint4*)&XX[4*idx]);
while (loop--)
switch(ALGO) {
case A_SCRYPT: xor_salsa8(b, c); xor_salsa8(c, b); break;
case A_SCRYPT_JANE: xor_chacha8(b, c); xor_chacha8(c, b); break;
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] ^= b[idx];
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] ^= c[idx];
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+Z])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+16+Z])) = *((ulonglong2*)XB[wu]);
}
template <int ALGO, int TEX_DIM> __global__ void
fermi_scrypt_core_kernelB_LG_tex(uint32_t *g_odata, unsigned int N, unsigned int LOOKUP_GAP)
{
extern __shared__ unsigned char x[];
uint32_t ((*X)[WU_PER_WARP][16+4]) = (uint32_t (*)[WU_PER_WARP][16+4]) x;
int warpIdx = threadIdx.x / warpSize;
int warpThread = threadIdx.x % warpSize;
// variables supporting the large memory transaction magic
unsigned int Y = warpThread/4;
unsigned int Z = 4*(warpThread%4);
// add block specific offsets
int WARPS_PER_BLOCK = blockDim.x / 32;
int offset = blockIdx.x * WU_PER_BLOCK + warpIdx * WU_PER_WARP;
g_odata += 32 * offset;
// registers to store an entire work unit
uint4 B[4], C[4];
uint32_t ((*XB)[16+4]) = (uint32_t (*)[16+4])&X[warpIdx][Y][Z];
uint32_t *XX = X[warpIdx][warpThread];
uint32_t pos = (N-1)/LOOKUP_GAP; uint32_t loop = 1 + (N-1)-pos*LOOKUP_GAP;
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + pos*32 + Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + pos*32 + 16+Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] = *((uint4*)&XX[4*idx]);
while (loop--)
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
for (int i = 0; i < N; i++) {
uint32_t j = C[0].x & (N-1);
uint32_t pos = j / LOOKUP_GAP; uint32_t loop = j - pos*LOOKUP_GAP;
XX[16] = 32 * pos;
uint4 b[4], c[4];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + XB[wu][16-Z] + Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) b[idx] = *((uint4*)&XX[4*idx]);
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8) { unsigned int loc = (SCRATCH*(offset+wu+Y) + XB[wu][16-Z] + 16+Z)/4;
*((uint4*)XB[wu]) = ((TEX_DIM == 1) ?
tex1Dfetch(texRef1D_4_V, loc) :
tex2D(texRef2D_4_V, 0.5f + (loc%TEXWIDTH), 0.5f + (loc/TEXWIDTH))); }
#pragma unroll 4
for (int idx=0; idx < 4; idx++) c[idx] = *((uint4*)&XX[4*idx]);
while (loop--)
switch(ALGO) {
case A_SCRYPT: xor_salsa8(b, c); xor_salsa8(c, b); break;
case A_SCRYPT_JANE: xor_chacha8(b, c); xor_chacha8(c, b); break;
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) B[idx] ^= b[idx];
#pragma unroll 4
for (int idx=0; idx < 4; idx++) C[idx] ^= c[idx];
switch(ALGO) {
case A_SCRYPT: xor_salsa8(B, C); xor_salsa8(C, B); break;
case A_SCRYPT_JANE: xor_chacha8(B, C); xor_chacha8(C, B); break;
}
}
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = B[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+Z])) = *((ulonglong2*)XB[wu]);
#pragma unroll 4
for (int idx=0; idx < 4; idx++) *((uint4*)&XX[4*idx]) = C[idx];
#pragma unroll 4
for (int wu=0; wu < 32; wu+=8)
*((ulonglong2*)(&g_odata[32*(wu+Y)+16+Z])) = *((ulonglong2*)XB[wu]);
}