Skip to content

Instantly share code, notes, and snippets.

View monkins1010's full-sized avatar

Chris monkins1010

View GitHub Profile
// VARIANT ALTERATIONS
#define VARIANT1_INIT(part) {}
#define VARIANT1_1(p) \
if (VARIANT > 0) { \
const uint8_t tmp1 = ((const uint8_t*)(p))[11]; \
static const uint32_t table1 = 0x86420; \
const uint8_t index1 = (((tmp1 >> 3) & 6) | (tmp1 & 1)) << 1; \
((uint8_t*)(p))[11] = tmp1 ^ ((table1 >> index1) & 0x30); \
}
template<size_t ITERATIONS, uint32_t MEM, uint32_t MASK, xmrig::Algo ALGO, uint8_t VARIANT>
#ifdef XMR_STAK_THREADS
__launch_bounds__( XMR_STAK_THREADS * 4 )
#endif
__global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b, uint32_t * d_ctx_state,
uint64_t startNonce, uint32_t * __restrict__ d_input )
{
__shared__ uint32_t sharedMemory[1024];
cn_aes_gpu_init( sharedMemory );
R"===(
/*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 3 of the License, or
* any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
@monkins1010
monkins1010 / hashtogpu.c
Created December 22, 2018 16:46
miner start
extern "C" int scanhash_verus(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done)
{
uint32_t _ALIGN(64) endiandata[35];
const char *pdata = (const char*)work->data;
uint32_t *ptarget = work->target;
int dev_id = device_map[thr_id];
unsigned char block_41970[] = { 0xfd, 0x40, 0x05, 0x01 };
uint8_t _ALIGN(64) full_data[140 + 3 + 1344] = { 0 };
uint8_t* sol_data = &full_data[140];
@monkins1010
monkins1010 / scanshashv2.1.c
Created January 2, 2019 07:19
scanhash routine for verus2, not working on multiple threads
static const int PROTOCOL_VERSION = 170002;
#include <cuda_helper.h>
#define EQNONCE_OFFSET 30 /* 27:34 */
#define NONCE_OFT EQNONCE_OFFSET
static bool init[MAX_GPUS] = { 0 };
static int valid_sols[MAX_GPUS] = { 0 };
static uint8_t _ALIGN(64) data_sols[MAX_GPUS][10][1536] = { 0 }; // 140+3+1344 required
#include <cuda_helper.h>
#define EQNONCE_OFFSET 30 /* 27:34 */
#define NONCE_OFT EQNONCE_OFFSET
static bool init[MAX_GPUS] = { 0 };
static u128 data_key[MAX_GPUS][VERUS_KEY_SIZE128] = { 0 }; // 552 required
extern "C" void GenNewCLKey(unsigned char *seedBytes32, u128 *keyback)
// mul2.cpp : Defines the entry point for the console application.
//
#include "stdafx.h"
// ConsoleApplication1.cpp : This file contains the 'main' function. Program execution begins and ends there.
//
#pragma warning (disable : 4146)
//#include "pch.h"
#include <iostream>
__device__ __forceinline__ uint64_t ROTL64(const uint64_t value, const int offset) {
uint2 result;
if (offset >= 32) {
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
}
else {
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
// kernel-interface: john SHA256d
#ifdef VECTORS4
typedef uint4 u;
#elif defined VECTORS2
typedef uint2 u;
#else
typedef uint u;
#endif
0x3da9BB5acdD969a4F37d989670859dD8502e95ee