sha256d CUDA implementation

本文详细介绍了一种基于CUDA的SHA256哈希算法实现,包括GPU并行计算的具体步骤,如数据处理、状态更新及非对称密钥操作等。通过分析,读者将了解到如何利用GPU加速SHA256运算,提升加密效率。

摘要生成于 C知道 ,由 DeepSeek-R1 满血版支持, 前往体验 >

/*
 * sha256d CUDA implementation.
 * tpruvot 2017
 */

#include <stdio.h>
#include <stdint.h>
#include <memory.h>

#include <cuda_helper.h>
#include <miner.h>

__constant__ static uint32_t __align__(8) c_midstate76[8];
__constant__ static uint32_t __align__(8) c_dataEnd80[4];

const __constant__  uint32_t __align__(8) c_H256[8] = {
    0x6A09E667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU,
    0x510E527FU, 0x9B05688CU, 0x1F83D9ABU, 0x5BE0CD19U
};
__constant__ static uint32_t __align__(8) c_K[64];
__constant__ static uint32_t __align__(8) c_target[2];
__device__ uint64_t d_target[1];

static uint32_t* d_resNonces[MAX_GPUS] = { 0 };

// ------------------------------------------------------------------------------------------------

static const uint32_t cpu_H256[8] = {
    0x6A09E667U, 0xBB67AE85U, 0x3C6EF372U, 0xA54FF53AU,
    0x510E527FU, 0x9B05688CU, 0x1F83D9ABU, 0x5BE0CD19U
};

static const uint32_t cpu_K[64] = {
    0x428A2F98, 0x71374491, 0xB5C0FBCF, 0xE9B5DBA5, 0x3956C25B, 0x59F111F1, 0x923F82A4, 0xAB1C5ED5,
    0xD807AA98, 0x12835B01, 0x243185BE, 0x550C7DC3, 0x72BE5D74, 0x80DEB1FE, 0x9BDC06A7, 0xC19BF174,
    0xE49B69C1, 0xEFBE4786, 0x0FC19DC6, 0x240CA1CC, 0x2DE92C6F, 0x4A7484AA, 0x5CB0A9DC, 0x76F988DA,
    0x983E5152, 0xA831C66D, 0xB00327C8, 0xBF597FC7, 0xC6E00BF3, 0xD5A79147, 0x06CA6351, 0x14292967,
    0x27B70A85, 0x2E1B2138, 0x4D2C6DFC, 0x53380D13, 0x650A7354, 0x766A0ABB, 0x81C2C92E, 0x92722C85,
    0xA2BFE8A1, 0xA81A664B, 0xC24B8B70, 0xC76C51A3, 0xD192E819, 0xD6990624, 0xF40E3585, 0x106AA070,
    0x19A4C116, 0x1E376C08, 0x2748774C, 0x34B0BCB5, 0x391C0CB3, 0x4ED8AA4A, 0x5B9CCA4F, 0x682E6FF3,
    0x748F82EE, 0x78A5636F, 0x84C87814, 0x8CC70208, 0x90BEFFFA, 0xA4506CEB, 0xBEF9A3F7, 0xC67178F2
};

#define ROTR ROTR32

__host__
static void sha256_step1_host(uint32_t a, uint32_t b, uint32_t c, uint32_t &d,
    uint32_t e, uint32_t f, uint32_t g, uint32_t &h,
    uint32_t in, const uint32_t Kshared)
{
    uint32_t t1,t2;
    uint32_t vxandx = (((f) ^ (g)) & (e)) ^ (g); // xandx(e, f, g);
    uint32_t bsg21 = ROTR(e, 6) ^ ROTR(e, 11) ^ ROTR(e, 25); // bsg2_1(e);
    uint32_t bsg20 = ROTR(a, 2) ^ ROTR(a, 13) ^ ROTR(a, 22); //bsg2_0(a);
    uint32_t andorv = ((b) & (c)) | (((b) | (c)) & (a)); //andor32(a,b,c);

    t1 = h + bsg21 + vxandx + Kshared + in;
    t2 = bsg20 + andorv;
    d = d + t1;
    h = t1 + t2;
}

__host__
static void sha256_step2_host(uint32_t a, uint32_t b, uint32_t c, uint32_t &d,
    uint32_t e, uint32_t f, uint32_t g, uint32_t &h,
    uint32_t* in, uint32_t pc, const uint32_t Kshared)
{
    uint32_t t1,t2;

    int pcidx1 = (pc-2)  & 0xF;
    int pcidx2 = (pc-7)  & 0xF;
    int pcidx3 = (pc-15) & 0xF;

    uint32_t inx0 = in[pc];
    uint32_t inx1 = in[pcidx1];
    uint32_t inx2 = in[pcidx2];
    uint32_t inx3 = in[pcidx3];

    uint32_t ssg21 = ROTR(inx1, 17) ^ ROTR(inx1, 19) ^ SPH_T32((inx1) >> 10); //ssg2_1(inx1);
    uint32_t ssg20 = ROTR(inx3, 7) ^ ROTR(inx3, 18) ^ SPH_T32((inx3) >> 3); //ssg2_0(inx3);
    uint32_t vxandx = (((f) ^ (g)) & (e)) ^ (g); // xandx(e, f, g);
    uint32_t bsg21 = ROTR(e, 6) ^ ROTR(e, 11) ^ ROTR(e, 25); // bsg2_1(e);
    uint32_t bsg20 = ROTR(a, 2) ^ ROTR(a, 13) ^ ROTR(a, 22); //bsg2_0(a);
    uint32_t andorv = ((b) & (c)) | (((b) | (c)) & (a)); //andor32(a,b,c);

    in[pc] = ssg21 + inx2 + ssg20 + inx0;

    t1 = h + bsg21 + vxandx + Kshared + in[pc];
    t2 = bsg20 + andorv;
    d =  d + t1;
    h = t1 + t2;
}

__host__
static void sha256_round_body_host(uint32_t* in, uint32_t* state, const uint32_t* Kshared)
{
    uint32_t a = state[0];
    uint32_t b = state[1];
    uint32_t c = state[2];
    uint32_t d = state[3];
    uint32_t e = state[4];
    uint32_t f = state[5];
    uint32_t g = state[6];
    uint32_t h = state[7];

    sha256_step1_host(a,b,c,d,e,f,g,h,in[ 0], Kshared[ 0]);
    sha256_step1_host(h,a,b,c,d,e,f,g,in[ 1], Kshared[ 1]);
    sha256_step1_host(g,h,a,b,c,d,e,f,in[ 2], Kshared[ 2]);
    sha256_step1_host(f,g,h,a,b,c,d,e,in[ 3], Kshared[ 3]);
    sha256_step1_host(e,f,g,h,a,b,c,d,in[ 4], Kshared[ 4]);
    sha256_step1_host(d,e,f,g,h,a,b,c,in[ 5], Kshared[ 5]);
    sha256_step1_host(c,d,e,f,g,h,a,b,in[ 6], Kshared[ 6]);
    sha256_step1_host(b,c,d,e,f,g,h,a,in[ 7], Kshared[ 7]);
    sha256_step1_host(a,b,c,d,e,f,g,h,in[ 8], Kshared[ 8]);
    sha256_step1_host(h,a,b,c,d,e,f,g,in[ 9], Kshared[ 9]);
    sha256_step1_host(g,h,a,b,c,d,e,f,in[10], Kshared[10]);
    sha256_step1_host(f,g,h,a,b,c,d,e,in[11], Kshared[11]);
    sha256_step1_host(e,f,g,h,a,b,c,d,in[12], Kshared[12]);
    sha256_step1_host(d,e,f,g,h,a,b,c,in[13], Kshared[13]);
    sha256_step1_host(c,d,e,f,g,h,a,b,in[14], Kshared[14]);
    sha256_step1_host(b,c,d,e,f,g,h,a,in[15], Kshared[15]);

    for (int i=0; i<3; i++)
    {
        sha256_step2_host(a,b,c,d,e,f,g,h,in,0, Kshared[16+16*i]);
        sha256_step2_host(h,a,b,c,d,e,f,g,in,1, Kshared[17+16*i]);
        sha256_step2_host(g,h,a,b,c,d,e,f,in,2, Kshared[18+16*i]);
        sha256_step2_host(f,g,h,a,b,c,d,e,in,3, Kshared[19+16*i]);
        sha256_step2_host(e,f,g,h,a,b,c,d,in,4, Kshared[20+16*i]);
        sha256_step2_host(d,e,f,g,h,a,b,c,in,5, Kshared[21+16*i]);
        sha256_step2_host(c,d,e,f,g,h,a,b,in,6, Kshared[22+16*i]);
        sha256_step2_host(b,c,d,e,f,g,h,a,in,7, Kshared[23+16*i]);
        sha256_step2_host(a,b,c,d,e,f,g,h,in,8, Kshared[24+16*i]);
        sha256_step2_host(h,a,b,c,d,e,f,g,in,9, Kshared[25+16*i]);
        sha256_step2_host(g,h,a,b,c,d,e,f,in,10,Kshared[26+16*i]);
        sha256_step2_host(f,g,h,a,b,c,d,e,in,11,Kshared[27+16*i]);
        sha256_step2_host(e,f,g,h,a,b,c,d,in,12,Kshared[28+16*i]);
        sha256_step2_host(d,e,f,g,h,a,b,c,in,13,Kshared[29+16*i]);
        sha256_step2_host(c,d,e,f,g,h,a,b,in,14,Kshared[30+16*i]);
        sha256_step2_host(b,c,d,e,f,g,h,a,in,15,Kshared[31+16*i]);
    }

    state[0] += a;
    state[1] += b;
    state[2] += c;
    state[3] += d;
    state[4] += e;
    state[5] += f;
    state[6] += g;
    state[7] += h;
}

#define xor3b(a,b,c) (a ^ b ^ c)

__device__ __forceinline__ uint32_t bsg2_0(const uint32_t x)
{
    return xor3b(ROTR32(x,2),ROTR32(x,13),ROTR32(x,22));
}

__device__ __forceinline__ uint32_t bsg2_1(const uint32_t x)
{
    return xor3b(ROTR32(x,6),ROTR32(x,11),ROTR32(x,25));
}

__device__ __forceinline__ uint32_t ssg2_0(const uint32_t x)
{
    return xor3b(ROTR32(x,7),ROTR32(x,18),(x>>3));
}

__device__ __forceinline__ uint32_t ssg2_1(const uint32_t x)
{
    return xor3b(ROTR32(x,17),ROTR32(x,19),(x>>10));
}

__device__ __forceinline__ uint32_t andor32(const uint32_t a, const uint32_t b, const uint32_t c)
{
    uint32_t result;
    asm("{\n\t"
        ".reg .u32 m,n,o;\n\t"
        "and.b32 m,  %1, %2;\n\t"
        " or.b32 n,  %1, %2;\n\t"
        "and.b32 o,   n, %3;\n\t"
        " or.b32 %0,  m, o ;\n\t"
        "}\n\t" : "=r"(result) : "r"(a), "r"(b), "r"(c)
    );
    return result;
}

__device__ __forceinline__ uint2 vectorizeswap(uint64_t v) {
    uint2 result;
    asm("mov.b64 {%0,%1},%2; \n\t"
        : "=r"(result.y), "=r"(result.x) : "l"(v));
    return result;
}

__device__
static void sha2_step1(uint32_t a, uint32_t b, uint32_t c, uint32_t &d, uint32_t e, uint32_t f, uint32_t g, uint32_t &h,
    uint32_t in, const uint32_t Kshared)
{
    uint32_t t1,t2;
    uint32_t vxandx = xandx(e, f, g);
    uint32_t bsg21 = bsg2_1(e);
    uint32_t bsg20 = bsg2_0(a);
    uint32_t andorv = andor32(a,b,c);

    t1 = h + bsg21 + vxandx + Kshared + in;
    t2 = bsg20 + andorv;
    d = d + t1;
    h = t1 + t2;
}

__device__
static void sha2_step2(uint32_t a, uint32_t b, uint32_t c, uint32_t &d, uint32_t e, uint32_t f, uint32_t g, uint32_t &h,
    uint32_t* in, uint32_t pc, const uint32_t Kshared)
{
    uint32_t t1,t2;

    int pcidx1 = (pc-2) & 0xF;
    int pcidx2 = (pc-7) & 0xF;
    int pcidx3 = (pc-15) & 0xF;

    uint32_t inx0 = in[pc];
    uint32_t inx1 = in[pcidx1];
    uint32_t inx2 = in[pcidx2];
    uint32_t inx3 = in[pcidx3];

    uint32_t ssg21 = ssg2_1(inx1);
    uint32_t ssg20 = ssg2_0(inx3);
    uint32_t vxandx = xandx(e, f, g);
    uint32_t bsg21 = bsg2_1(e);
    uint32_t bsg20 = bsg2_0(a);
    uint32_t andorv = andor32(a,b,c);

    in[pc] = ssg21 + inx2 + ssg20 + inx0;

    t1 = h + bsg21 + vxandx + Kshared + in[pc];
    t2 = bsg20 + andorv;
    d =  d + t1;
    h = t1 + t2;
}

__device__
static void sha256_round_body(uint32_t* in, uint32_t* state, uint32_t* const Kshared)
{
    uint32_t a = state[0];
    uint32_t b = state[1];
    uint32_t c = state[2];
    uint32_t d = state[3];
    uint32_t e = state[4];
    uint32_t f = state[5];
    uint32_t g = state[6];
    uint32_t h = state[7];

    sha2_step1(a,b,c,d,e,f,g,h,in[ 0], Kshared[ 0]);
    sha2_step1(h,a,b,c,d,e,f,g,in[ 1], Kshared[ 1]);
    sha2_step1(g,h,a,b,c,d,e,f,in[ 2], Kshared[ 2]);
    sha2_step1(f,g,h,a,b,c,d,e,in[ 3], Kshared[ 3]);
    sha2_step1(e,f,g,h,a,b,c,d,in[ 4], Kshared[ 4]);
    sha2_step1(d,e,f,g,h,a,b,c,in[ 5], Kshared[ 5]);
    sha2_step1(c,d,e,f,g,h,a,b,in[ 6], Kshared[ 6]);
    sha2_step1(b,c,d,e,f,g,h,a,in[ 7], Kshared[ 7]);
    sha2_step1(a,b,c,d,e,f,g,h,in[ 8], Kshared[ 8]);
    sha2_step1(h,a,b,c,d,e,f,g,in[ 9], Kshared[ 9]);
    sha2_step1(g,h,a,b,c,d,e,f,in[10], Kshared[10]);
    sha2_step1(f,g,h,a,b,c,d,e,in[11], Kshared[11]);
    sha2_step1(e,f,g,h,a,b,c,d,in[12], Kshared[12]);
    sha2_step1(d,e,f,g,h,a,b,c,in[13], Kshared[13]);
    sha2_step1(c,d,e,f,g,h,a,b,in[14], Kshared[14]);
    sha2_step1(b,c,d,e,f,g,h,a,in[15], Kshared[15]);

    #pragma unroll
    for (int i=0; i<3; i++)
    {
        sha2_step2(a,b,c,d,e,f,g,h,in,0, Kshared[16+16*i]);
        sha2_step2(h,a,b,c,d,e,f,g,in,1, Kshared[17+16*i]);
        sha2_step2(g,h,a,b,c,d,e,f,in,2, Kshared[18+16*i]);
        sha2_step2(f,g,h,a,b,c,d,e,in,3, Kshared[19+16*i]);
        sha2_step2(e,f,g,h,a,b,c,d,in,4, Kshared[20+16*i]);
        sha2_step2(d,e,f,g,h,a,b,c,in,5, Kshared[21+16*i]);
        sha2_step2(c,d,e,f,g,h,a,b,in,6, Kshared[22+16*i]);
        sha2_step2(b,c,d,e,f,g,h,a,in,7, Kshared[23+16*i]);
        sha2_step2(a,b,c,d,e,f,g,h,in,8, Kshared[24+16*i]);
        sha2_step2(h,a,b,c,d,e,f,g,in,9, Kshared[25+16*i]);
        sha2_step2(g,h,a,b,c,d,e,f,in,10,Kshared[26+16*i]);
        sha2_step2(f,g,h,a,b,c,d,e,in,11,Kshared[27+16*i]);
        sha2_step2(e,f,g,h,a,b,c,d,in,12,Kshared[28+16*i]);
        sha2_step2(d,e,f,g,h,a,b,c,in,13,Kshared[29+16*i]);
        sha2_step2(c,d,e,f,g,h,a,b,in,14,Kshared[30+16*i]);
        sha2_step2(b,c,d,e,f,g,h,a,in,15,Kshared[31+16*i]);
    }

    state[0] += a;
    state[1] += b;
    state[2] += c;
    state[3] += d;
    state[4] += e;
    state[5] += f;
    state[6] += g;
    state[7] += h;
}

__device__
static void sha256_round_last(uint32_t* in, uint32_t* state, uint32_t* const Kshared)
{
    uint32_t a = state[0];
    uint32_t b = state[1];
    uint32_t c = state[2];
    uint32_t d = state[3];
    uint32_t e = state[4];
    uint32_t f = state[5];
    uint32_t g = state[6];
    uint32_t h = state[7];

    sha2_step1(a,b,c,d, e,f,g,h, in[ 0], Kshared[ 0]);
    sha2_step1(h,a,b,c, d,e,f,g, in[ 1], Kshared[ 1]);
    sha2_step1(g,h,a,b, c,d,e,f, in[ 2], Kshared[ 2]);
    sha2_step1(f,g,h,a, b,c,d,e, in[ 3], Kshared[ 3]);
    sha2_step1(e,f,g,h, a,b,c,d, in[ 4], Kshared[ 4]);
    sha2_step1(d,e,f,g, h,a,b,c, in[ 5], Kshared[ 5]);
    sha2_step1(c,d,e,f, g,h,a,b, in[ 6], Kshared[ 6]);
    sha2_step1(b,c,d,e, f,g,h,a, in[ 7], Kshared[ 7]);
    sha2_step1(a,b,c,d, e,f,g,h, in[ 8], Kshared[ 8]);
    sha2_step1(h,a,b,c, d,e,f,g, in[ 9], Kshared[ 9]);
    sha2_step1(g,h,a,b, c,d,e,f, in[10], Kshared[10]);
    sha2_step1(f,g,h,a, b,c,d,e, in[11], Kshared[11]);
    sha2_step1(e,f,g,h, a,b,c,d, in[12], Kshared[12]);
    sha2_step1(d,e,f,g, h,a,b,c, in[13], Kshared[13]);
    sha2_step1(c,d,e,f, g,h,a,b, in[14], Kshared[14]);
    sha2_step1(b,c,d,e, f,g,h,a, in[15], Kshared[15]);

    #pragma unroll
    for (int i=0; i<2; i++)
    {
        sha2_step2(a,b,c,d, e,f,g,h, in, 0, Kshared[16+16*i]);
        sha2_step2(h,a,b,c, d,e,f,g, in, 1, Kshared[17+16*i]);
        sha2_step2(g,h,a,b, c,d,e,f, in, 2, Kshared[18+16*i]);
        sha2_step2(f,g,h,a, b,c,d,e, in, 3, Kshared[19+16*i]);
        sha2_step2(e,f,g,h, a,b,c,d, in, 4, Kshared[20+16*i]);
        sha2_step2(d,e,f,g, h,a,b,c, in, 5, Kshared[21+16*i]);
        sha2_step2(c,d,e,f, g,h,a,b, in, 6, Kshared[22+16*i]);
        sha2_step2(b,c,d,e, f,g,h,a, in, 7, Kshared[23+16*i]);
        sha2_step2(a,b,c,d, e,f,g,h, in, 8, Kshared[24+16*i]);
        sha2_step2(h,a,b,c, d,e,f,g, in, 9, Kshared[25+16*i]);
        sha2_step2(g,h,a,b, c,d,e,f, in,10, Kshared[26+16*i]);
        sha2_step2(f,g,h,a, b,c,d,e, in,11, Kshared[27+16*i]);
        sha2_step2(e,f,g,h, a,b,c,d, in,12, Kshared[28+16*i]);
        sha2_step2(d,e,f,g, h,a,b,c, in,13, Kshared[29+16*i]);
        sha2_step2(c,d,e,f, g,h,a,b, in,14, Kshared[30+16*i]);
        sha2_step2(b,c,d,e, f,g,h,a, in,15, Kshared[31+16*i]);
    }

    sha2_step2(a,b,c,d, e,f,g,h, in, 0, Kshared[16+16*2]);
    sha2_step2(h,a,b,c, d,e,f,g, in, 1, Kshared[17+16*2]);
    sha2_step2(g,h,a,b, c,d,e,f, in, 2, Kshared[18+16*2]);
    sha2_step2(f,g,h,a, b,c,d,e, in, 3, Kshared[19+16*2]);
    sha2_step2(e,f,g,h, a,b,c,d, in, 4, Kshared[20+16*2]);
    sha2_step2(d,e,f,g, h,a,b,c, in, 5, Kshared[21+16*2]);
    sha2_step2(c,d,e,f, g,h,a,b, in, 6, Kshared[22+16*2]);
    sha2_step2(b,c,d,e, f,g,h,a, in, 7, Kshared[23+16*2]);
    sha2_step2(a,b,c,d, e,f,g,h, in, 8, Kshared[24+16*2]);
    sha2_step2(h,a,b,c, d,e,f,g, in, 9, Kshared[25+16*2]);
    sha2_step2(g,h,a,b, c,d,e,f, in,10, Kshared[26+16*2]);
    sha2_step2(f,g,h,a, b,c,d,e, in,11, Kshared[27+16*2]);
    sha2_step2(e,f,g,h, a,b,c,d, in,12, Kshared[28+16*2]);
    sha2_step2(d,e,f,g, h,a,b,c, in,13, Kshared[29+16*2]);

    state[6] += g;
    state[7] += h;
}

__device__ __forceinline__
uint64_t cuda_swab32ll(uint64_t x) {
    return MAKE_ULONGLONG(cuda_swab32(_LODWORD(x)), cuda_swab32(_HIDWORD(x)));
}

__global__
/*__launch_bounds__(256,3)*/
void sha256d_gpu_hash_shared(const uint32_t threads, const uint32_t startNonce, uint32_t *resNonces)
{

    const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);

    __shared__ uint32_t s_K[64*4];
    //s_K[thread & 63] = c_K[thread & 63];
    if (threadIdx.x < 64U) s_K[threadIdx.x] = c_K[threadIdx.x];

    if (thread < threads)
    {
        const uint32_t nonce = startNonce + thread;

        uint32_t dat[16];
        AS_UINT2(dat) = AS_UINT2(c_dataEnd80);
        dat[ 2] = c_dataEnd80[2];
        dat[ 3] = nonce;
        dat[ 4] = 0x80000000;
        dat[15] = 0x280;
        #pragma unroll
        for (int i=5; i<15; i++) dat[i] = 0;

        uint32_t buf[8];
        #pragma unroll
        for (int i=0; i<8; i+=2) AS_UINT2(&buf[i]) = AS_UINT2(&c_midstate76[i]);
        //for (int i=0; i<8; i++) buf[i] = c_midstate76[i];

        sha256_round_body(dat, buf, s_K);

        // second sha256

        #pragma unroll
        for (int i=0; i<8; i++) dat[i] = buf[i];
        dat[8] = 0x80000000;
        #pragma unroll
        for (int i=9; i<15; i++) dat[i] = 0;
        dat[15] = 0x100;

        #pragma unroll
        for (int i=0; i<8; i++) buf[i] = c_H256[i];

        sha256_round_last(dat, buf, s_K);

        // valid nonces
        uint64_t high = cuda_swab32ll(((uint64_t*)buf)[3]);
        if (high <= c_target[0]) {
            //printf("%08x %08x - %016llx %016llx - %08x %08x\n", buf[7], buf[6], high, d_target[0], c_target[1], c_target[0]);
            resNonces[1] = atomicExch(resNonces, nonce);
            //d_target[0] = high;
        }
    }

}

__host__
void sha256d_init(int thr_id)
{
    cuda_get_arch(thr_id);
    cudaMemcpyToSymbol(c_K, cpu_K, sizeof(cpu_K), 0, cudaMemcpyHostToDevice);
    CUDA_SAFE_CALL(cudaMalloc(&d_resNonces[thr_id], 2*sizeof(uint32_t)));
}

__host__
void sha256d_free(int thr_id)
{
    if (d_resNonces[thr_id]) cudaFree(d_resNonces[thr_id]);
    d_resNonces[thr_id] = NULL;
}

__host__
void sha256d_setBlock_80(uint32_t *pdata, uint32_t *ptarget)
{
    uint32_t _ALIGN(64) in[16], buf[8], end[4];
    for (int i=0;i<16;i++) in[i] = cuda_swab32(pdata[i]);
    for (int i=0;i<8;i++) buf[i] = cpu_H256[i];
    for (int i=0;i<4;i++) end[i] = cuda_swab32(pdata[16+i]);
    sha256_round_body_host(in, buf, cpu_K);

    CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_midstate76, buf, 32, 0, cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_dataEnd80,  end, sizeof(end), 0, cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_target, &ptarget[6], 8, 0, cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_target, &ptarget[6], 8, 0, cudaMemcpyHostToDevice));
}

__host__
void sha256d_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *resNonces)
{
    const uint32_t threadsperblock = 256;

    dim3 grid(threads/threadsperblock);
    dim3 block(threadsperblock);

    CUDA_SAFE_CALL(cudaMemset(d_resNonces[thr_id], 0xFF, 2 * sizeof(uint32_t)));
    cudaThreadSynchronize();
    sha256d_gpu_hash_shared <<<grid, block>>> (threads, startNonce, d_resNonces[thr_id]);
    cudaThreadSynchronize();

    CUDA_SAFE_CALL(cudaMemcpy(resNonces, d_resNonces[thr_id], 2 * sizeof(uint32_t), cudaMemcpyDeviceToHost));
    if (resNonces[0] == resNonces[1]) {
        resNonces[1] = UINT32_MAX;
    }
}
 

(unsloth) root@autodl-container-578c43a3a7-817faf39:~/autodl-tmp/llm_train/src# vllm serve ./ckpts/qwen3-0.6b --enable-auto-tool-choice --tool-call-parser hermes INFO 07-31 23:06:38 [__init__.py:235] Automatically detected platform cuda. INFO 07-31 23:06:41 [api_server.py:1755] vLLM API server version 0.10.0 INFO 07-31 23:06:41 [cli_args.py:261] non-default args: {'model_tag': './ckpts/qwen3-0.6b', 'enable_auto_tool_choice': True, 'tool_call_parser': 'hermes', 'model': './ckpts/qwen3-0.6b'} INFO 07-31 23:06:49 [config.py:1604] Using max model len 40960 WARNING 07-31 23:06:49 [config.py:1084] bitsandbytes quantization is not fully optimized yet. The speed can be slower than non-quantized models. INFO 07-31 23:06:49 [config.py:2434] Chunked prefill is enabled with max_num_batched_tokens=2048. INFO 07-31 23:06:54 [__init__.py:235] Automatically detected platform cuda. INFO 07-31 23:06:57 [core.py:572] Waiting for init message from front-end. INFO 07-31 23:06:57 [core.py:71] Initializing a V1 LLM engine (v0.10.0) with config: model='./ckpts/qwen3-0.6b', speculative_config=None, tokenizer='./ckpts/qwen3-0.6b', skip_tokenizer_init=False, tokenizer_mode=auto, revision=None, override_neuron_config={}, tokenizer_revision=None, trust_remote_code=False, dtype=torch.bfloat16, max_seq_len=40960, download_dir=None, load_format=LoadFormat.BITSANDBYTES, tensor_parallel_size=1, pipeline_parallel_size=1, disable_custom_all_reduce=False, quantization=bitsandbytes, enforce_eager=False, kv_cache_dtype=auto, device_config=cuda, decoding_config=DecodingConfig(backend='auto', disable_fallback=False, disable_any_whitespace=False, disable_additional_properties=False, reasoning_backend=''), observability_config=ObservabilityConfig(show_hidden_metrics_for_version=None, otlp_traces_endpoint=None, collect_detailed_traces=None), seed=0, served_model_name=./ckpts/qwen3-0.6b, num_scheduler_steps=1, multi_step_stream_outputs=True, enable_prefix_caching=True, chunked_prefill_enabled=True, use_async_output_proc=True, pooler_config=None, compilation_config={"level":3,"debug_dump_path":"","cache_dir":"","backend":"","custom_ops":[],"splitting_ops":["vllm.unified_attention","vllm.unified_attention_with_output","vllm.mamba_mixer2"],"use_inductor":true,"compile_sizes":[],"inductor_compile_config":{"enable_auto_functionalized_v2":false},"inductor_passes":{},"use_cudagraph":true,"cudagraph_num_of_warmups":1,"cudagraph_capture_sizes":[512,504,496,488,480,472,464,456,448,440,432,424,416,408,400,392,384,376,368,360,352,344,336,328,320,312,304,296,288,280,272,264,256,248,240,232,224,216,208,200,192,184,176,168,160,152,144,136,128,120,112,104,96,88,80,72,64,56,48,40,32,24,16,8,4,2,1],"cudagraph_copy_inputs":false,"full_cuda_graph":false,"max_capture_size":512,"local_cache_dir":null} INFO 07-31 23:06:58 [parallel_state.py:1102] rank 0 in world size 1 is assigned as DP rank 0, PP rank 0, TP rank 0, EP rank 0 WARNING 07-31 23:06:58 [topk_topp_sampler.py:59] FlashInfer is not available. Falling back to the PyTorch-native implementation of top-p & top-k sampling. For the best performance, please install FlashInfer. INFO 07-31 23:06:58 [gpu_model_runner.py:1843] Starting to load model ./ckpts/qwen3-0.6b... INFO 07-31 23:06:58 [gpu_model_runner.py:1875] Loading model from scratch... INFO 07-31 23:06:58 [cuda.py:290] Using Flash Attention backend on V1 engine. INFO 07-31 23:06:58 [bitsandbytes_loader.py:733] Loading weights with BitsAndBytes quantization. May take a while ... Loading safetensors checkpoint shards: 0% Completed | 0/1 [00:00<?, ?it/s] Loading safetensors checkpoint shards: 100% Completed | 1/1 [00:00<00:00, 31.45it/s] Loading safetensors checkpoint shards: 0% Completed | 0/1 [00:00<?, ?it/s] ERROR 07-31 23:06:59 [core.py:632] EngineCore failed to start. ERROR 07-31 23:06:59 [core.py:632] Traceback (most recent call last): ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 623, in run_engine_core ERROR 07-31 23:06:59 [core.py:632] engine_core = EngineCoreProc(*args, **kwargs) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 441, in __init__ ERROR 07-31 23:06:59 [core.py:632] super().__init__(vllm_config, executor_class, log_stats, ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 77, in __init__ ERROR 07-31 23:06:59 [core.py:632] self.model_executor = executor_class(vllm_config) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/executor_base.py", line 53, in __init__ ERROR 07-31 23:06:59 [core.py:632] self._init_executor() ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/uniproc_executor.py", line 49, in _init_executor ERROR 07-31 23:06:59 [core.py:632] self.collective_rpc("load_model") ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/uniproc_executor.py", line 58, in collective_rpc ERROR 07-31 23:06:59 [core.py:632] answer = run_method(self.driver_worker, method, args, kwargs) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/utils/__init__.py", line 2985, in run_method ERROR 07-31 23:06:59 [core.py:632] return func(*args, **kwargs) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/worker/gpu_worker.py", line 201, in load_model ERROR 07-31 23:06:59 [core.py:632] self.model_runner.load_model(eep_scale_up=eep_scale_up) ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/worker/gpu_model_runner.py", line 1876, in load_model ERROR 07-31 23:06:59 [core.py:632] self.model = model_loader.load_model( ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/model_loader/base_loader.py", line 49, in load_model ERROR 07-31 23:06:59 [core.py:632] self.load_weights(model, model_config) ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/model_loader/bitsandbytes_loader.py", line 741, in load_weights ERROR 07-31 23:06:59 [core.py:632] loaded_weights = model.load_weights(qweight_iterator) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/qwen3.py", line 321, in load_weights ERROR 07-31 23:06:59 [core.py:632] return loader.load_weights(weights) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 291, in load_weights ERROR 07-31 23:06:59 [core.py:632] autoloaded_weights = set(self._load_module("", self.module, weights)) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 249, in _load_module ERROR 07-31 23:06:59 [core.py:632] yield from self._load_module(prefix, ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 222, in _load_module ERROR 07-31 23:06:59 [core.py:632] loaded_params = module_load_weights(weights) ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/qwen2.py", line 404, in load_weights ERROR 07-31 23:06:59 [core.py:632] weight_loader(param, loaded_weight, shard_id) ERROR 07-31 23:06:59 [core.py:632] File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/layers/linear.py", line 806, in weight_loader ERROR 07-31 23:06:59 [core.py:632] assert param_data.shape == loaded_weight.shape ERROR 07-31 23:06:59 [core.py:632] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ERROR 07-31 23:06:59 [core.py:632] AssertionError Process EngineCore_0: Traceback (most recent call last): File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/multiprocessing/process.py", line 314, in _bootstrap self.run() File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/multiprocessing/process.py", line 108, in run self._target(*self._args, **self._kwargs) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 636, in run_engine_core raise e File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 623, in run_engine_core engine_core = EngineCoreProc(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 441, in __init__ super().__init__(vllm_config, executor_class, log_stats, File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core.py", line 77, in __init__ self.model_executor = executor_class(vllm_config) ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/executor_base.py", line 53, in __init__ self._init_executor() File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/uniproc_executor.py", line 49, in _init_executor self.collective_rpc("load_model") File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/executor/uniproc_executor.py", line 58, in collective_rpc answer = run_method(self.driver_worker, method, args, kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/utils/__init__.py", line 2985, in run_method return func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/worker/gpu_worker.py", line 201, in load_model self.model_runner.load_model(eep_scale_up=eep_scale_up) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/worker/gpu_model_runner.py", line 1876, in load_model self.model = model_loader.load_model( ^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/model_loader/base_loader.py", line 49, in load_model self.load_weights(model, model_config) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/model_loader/bitsandbytes_loader.py", line 741, in load_weights loaded_weights = model.load_weights(qweight_iterator) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/qwen3.py", line 321, in load_weights return loader.load_weights(weights) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 291, in load_weights autoloaded_weights = set(self._load_module("", self.module, weights)) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 249, in _load_module yield from self._load_module(prefix, File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/utils.py", line 222, in _load_module loaded_params = module_load_weights(weights) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/models/qwen2.py", line 404, in load_weights weight_loader(param, loaded_weight, shard_id) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/model_executor/layers/linear.py", line 806, in weight_loader assert param_data.shape == loaded_weight.shape ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ AssertionError Loading safetensors checkpoint shards: 0% Completed | 0/1 [00:00<?, ?it/s] [rank0]:[W731 23:07:00.371522631 ProcessGroupNCCL.cpp:1479] Warning: WARNING: destroy_process_group() was not called before program exit, which can leak resources. For more info, please see https://pytorch.org/docs/stable/distributed.html#shutdown (function operator()) Traceback (most recent call last): File "/root/autodl-tmp/miniconda3/envs/unsloth/bin/vllm", line 8, in <module> sys.exit(main()) ^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/cli/main.py", line 54, in main args.dispatch_function(args) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/cli/serve.py", line 52, in cmd uvloop.run(run_server(args)) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/uvloop/__init__.py", line 105, in run return runner.run(wrapper()) ^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/asyncio/runners.py", line 118, in run return self._loop.run_until_complete(task) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "uvloop/loop.pyx", line 1518, in uvloop.loop.Loop.run_until_complete File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/uvloop/__init__.py", line 61, in wrapper return await main ^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/openai/api_server.py", line 1791, in run_server await run_server_worker(listen_address, sock, args, **uvicorn_kwargs) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/openai/api_server.py", line 1811, in run_server_worker async with build_async_engine_client(args, client_config) as engine_client: File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/contextlib.py", line 210, in __aenter__ return await anext(self.gen) ^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/openai/api_server.py", line 158, in build_async_engine_client async with build_async_engine_client_from_engine_args( File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/contextlib.py", line 210, in __aenter__ return await anext(self.gen) ^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/entrypoints/openai/api_server.py", line 194, in build_async_engine_client_from_engine_args async_llm = AsyncLLM.from_vllm_config( ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/async_llm.py", line 163, in from_vllm_config return cls( ^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/async_llm.py", line 117, in __init__ self.engine_core = EngineCoreClient.make_async_mp_client( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core_client.py", line 98, in make_async_mp_client return AsyncMPClient(*client_args) ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core_client.py", line 677, in __init__ super().__init__( File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/core_client.py", line 408, in __init__ with launch_core_engines(vllm_config, executor_class, File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/contextlib.py", line 144, in __exit__ next(self.gen) File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/utils.py", line 697, in launch_core_engines wait_for_engine_startup( File "/root/autodl-tmp/miniconda3/envs/unsloth/lib/python3.11/site-packages/vllm/v1/engine/utils.py", line 750, in wait_for_engine_startup raise RuntimeError("Engine core initialization failed. " RuntimeError: Engine core initialization failed. See root cause above. Failed core proc(s): {}
最新发布
08-01
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值