[原创]一篇关于TencentPediyKeygenMe的贴子-CTF对抗-看雪安全社区

[原创]一篇关于TencentPediyKeygenMe的贴子

发表于: 2026-5-6 11:57 9629

[原创]一篇关于TencentPediyKeygenMe的贴子

n00bzx

2026-5-6 11:57

9629

好耶!

花了半个月时间,把这题出了.但是毕竟题还是要自己做嘛,不能给脚本和flag.那就失去了学习的必要了.做题毕竟还是为了学习嘛,对吗?

先说这题.其他算法不必多说,文章将聚焦于其中素数域椭圆曲线的攻击,不提供题目相关思路和flag.

本文的代码在语法上不保证能够100%正确运行,所以请各位喜欢ctrl-c+ctrl-v的大师傅们用之前记得改一改啊...#(滑稽)

下文为方便,pollard-rho算法和pollard-kangaroo算法均简称为rho和kangaroo.

听说,有个工具叫sagemath.听说,他是用python写的.听说,核心代码是09年的.#(滑稽)

于是,我开始了逆向...

不行了,他的python代码一大堆注释,而且代码结构松散...老开源都这样...

这是修改后的rho(少量优化):

```
seed=0x1145141919810666
def my_rand(a,b):
    global seed
    seed=(seed*6364136223846793005+1)&0xffffffffffffffff
    return seed%(b-a)+a
def discrete_log_rho(pub, G):
    from sage.rings.finite_rings.integer_mod_ring import IntegerModRing
    partition_size = 20
    memory_size = 4
    n = G.order()
    if n.is_prime()!=True:
        raise ValueError("no!")
    reset_bound = 8 * n.isqrt()
    I = IntegerModRing(n)
    for _ in range(10):
        arr1 = [my_rand(0,n) for i in range(partition_size)]
        arr2 = [my_rand(0,n) for i in range(partition_size)]
        M = [G*arr1[i]+pub*arr2[i] for i in range(partition_size)]
        ax = my_rand(0,n)
        x = G*ax
        bx = 0
        sigma = [(0, None)] * memory_size
        H = {}
        i0 = 0
        nextsigma = 0
        for i in range(reset_bound):
            combined = int(x.xy()[0]) * n + int(x.xy()[1])
            s = combined % partition_size
            x+=M[s]
            ax=(ax + arr1[s])%n
            bx=(bx + arr2[s])%n
            if x in H:
                ay, by = H[x]
                if bx == by:
                    break
                else:
                    res = sage.rings.integer.Integer(I((ay - ax)%n) / I((bx - by)%n))
                    if G*res == pub:
                        return res
                    else:
                        break
            elif i >= nextsigma:
                if sigma[i0][1] != None:
                    H.pop(sigma[i0][1])
                sigma[i0] = (i, x)
                i0 = (i0 + 1) % memory_size
                nextsigma = (3 * sigma[i0][0])%n
                H[x] = (ax, bx)
    print("no!")
```

这是修改后的kangaroo(又叫lambda)(少量优化):

```
def zx_hash(x):
    pos=x.xy()
    return int(pos[0])^int(pos[1])
def discrete_log_lambda(a,n, base, bounds):
    from sage.rings.integer import Integer
    from operator import mul, add
    lb, ub = bounds
    width = Integer(ub - lb)
    N = width.isqrt() + 1
    for s in range(8):
        M = []
        k=0
        while (1<<k)<N:
            r=my_rand(1,N)
            M.append((r, mul(base, r)))
            k += 1
        H = mul(base, ub)
        c = ub
        for i in range(N):
            r, e = M[zx_hash(H) % k]
            print(H,k)
            H = add(H, e)
            c += r
        mem = H
        H = a
        d = 0
        while True:
            r, e = M[zx_hash(H) % k]
            H = add(H, e)
            d += r
            res = c - d
            if res < lb:
                break
            if ub > res and H == mem:
                return res
    print("no!")
```

自动生成曲线测试下rho的速度...

```
import random
while True:
    p = random_prime(2^32)#随机有限域阶p
    a = random.randrange(-pow(2,63),pow(2,63))#模p之后在必定在8字节内
    b = random.randrange(-pow(2,63),pow(2,63))
    E = EllipticCurve(GF(p), [a,b])#曲线
    G = E.gens()[0]#基点
    n = G.order()#阶
    priv = random.randrange(n)#随机私钥
    pub = priv * G#对应公钥
    pub_x=int(pub.xy()[0])#公钥x坐标
    pub_y=int(pub.xy()[1])#公钥y坐标
    G_x=int(G.xy()[0])#基点x坐标
    G_y=int(G.xy()[1])#基点y坐标
    if n.is_prime():#为素数域
        print(G_x,G_y,pub_x,pub_y,a%int(p),int(p),int(n),priv,b)#输出信息
        print(discrete_log_rho(pub,G)==priv)#攻击并验证
```

好慢啊...

rho跑2^32域的速度大约为2秒一个,勉强能跑...

kangaroo范围大一点就出不来了...就不贴出来了...

不行不行...这怎么做题,题目的域可是2^64级的...

转成c先,部分算法采用汇编实现...

以下代码全部最大域限制为2^64...

整数除法,128位除以64位...
```
_div proc
    push rbx
    xor rdi,rdi
    xor rax,rax
    xor r9,r9
    mov r10,80h
lp:
    shl rdi,1
    setc bl
    shl rdx,1
    rcl rcx,1
    adc rdi,0
    mov rsi,rdi
    sub rdi,r8
    setnc bh
    or bl,bh
    test bl,bl
    jnz na
    add rdi,r8
na:
    shl rax,1
    rcl r9,1
    or al,bl
    dec r10
    jnz lp
    movq xmm0,rax
    movq xmm1,r9
    punpcklqdq xmm0,xmm1
    pop rbx
    ret
_div endp
```
整数平方根...
```
_sqrt proc;龟速版,rdi,r10,r11
    xor rax,rax
    mov r10,80000000h
    mov r11,rcx
    mov rcx,1fh
lp:
    mov rdi,rax
    shl rdi,1
    add rdi,r10
    shl rdi,cl
    cmp r11,rdi
    jc b
    add rax,r10
    sub r11,rdi
b:
    shr r10,1
    dec rcx
    test r10,r10
    jnz lp
    ret
_sqrt endp
```
齐次坐标转仿射坐标...
```
point_to_affine proc
    mov rcx,[r15+350h]
    mov rdx,rcx
    call mod_mul_p
    mov r8,rax
    mov rcx,r8
    call pow_mod_p
    mov rcx,[r15+340h]
    mov rdx,rax
    call mod_mul_p
    mov [r15+18h],rax
    mov rcx,r8
    mov rdx,[r15+350h]
    call mod_mul_p
    mov rcx,rax
    call pow_mod_p
    mov rcx,[r15+348h]
    mov rdx,rax
    call mod_mul_p
    mov [r15+20h],rax
    ret
point_to_affine endp
```
倍点...
```
pd proc
    push rsi
    push rdi
    push r12
    push r13
    push r14
    mov rcx,[r15+348h]
    mov rdx,2
    call mod_mul_p

mov rsi,rax
    mov rcx,rsi
    mov rdx,[r15+350h]
    call mod_mul_p

mov [r15+38h],rax
    mov rcx,[r15+350h]
    mov rdx,[r15+350h]
    call mod_mul_p

mov rsi,rax
    mov rcx,rsi
    mov rdx,rsi
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,A
    call mod_mul_p

mov r12,rax
    mov rcx,[r15+340h]
    mov rdx,[r15+340h]
    call mod_mul_p

mov rdi,rax
    mov rcx,rdi
    mov rdx,3
    call mod_mul_p

mov rdi,rax
    mov rcx,rdi
    mov rdx,r12
    call mod_add_p

mov rsi,rax
    mov rcx,[r15+348h]
    mov rdx,[r15+348h]
    call mod_mul_p

mov rdi,rax
    mov rcx,[r15+340h]
    mov rdx,4
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,rdi
    call mod_mul_p

mov r12,rax
    mov rcx,rdi
    mov rdx,rdi
    call mod_mul_p

mov r14,rax
    mov rcx,r14
    mov rdx,8
    call mod_mul_p

mov r14,rax
    mov rcx,r12
    mov rdx,2
    call mod_mul_p

mov rdi,rax
    mov rcx,rsi
    mov rdx,rsi
    call mod_mul_p

mov r13,rax
    mov rcx,r13
    mov rdx,rdi
    call mod_sub_p

mov [r15+28h],rax
    mov rcx,r12
    mov rdx,[r15+28h]
    call mod_sub_p

mov r13,rax
    mov rcx,r13
    mov rdx,rsi
    call mod_mul_p

mov r13,rax
    mov rcx,r13
    mov rdx,r14
    call mod_sub_p

mov [r15+30h],rax
    pop r14
    pop r13
    pop r12
    pop rdi
    pop rsi
    ret
pd endp
```
倍点对应的c代码(不完全对应)...
```
__declspec(noinline) void pd()
{
    DWORD64 y2 = mod_mul_p(ip31_y, ip31_y);
    DWORD64 z2 = mod_mul_p(ip31_z, ip31_z);
    DWORD64 L1 = mod_add_p(mod_mul_p(3, mod_mul_p(ip31_x, ip31_x)), mod_mul_p(A, mod_mul_p(z2, z2)));
    DWORD64 L2 = mod_mul_p(mod_mul_p(4, ip31_x), y2);
    DWORD64 L3 = mod_mul_p(8, mod_mul_p(y2, y2));
    rp3_x = mod_sub_p(mod_mul_p(L1, L1), mod_mul_p(2, L2));
    rp3_y = mod_sub_p(mod_mul_p(L1, mod_sub_p(L2, rp3_x)), L3);
    rp3_z = mod_mul_p(mod_mul_p(2, ip31_y), ip31_z);
}
```
点加...
```
pa proc
    push rsi
    push rdi
    push rbp
    push r12
    push r13
    push r14

mov rcx,[r15+368h]
    mov rdx,[r15+368h]
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,[r15+340h]
    call mod_mul_p

mov rdi,rax
    mov rdx,[r15+348h]
    call mod_mul_p

mov rbp,rax
    mov rcx,rbp
    mov rdx,[r15+368h]
    call mod_mul_p

mov rbp,rax
    mov rcx,[r15+350h]
    mov rdx,[r15+350h]
    call mod_mul_p

mov rsi,rax
    mov rcx,rsi
    mov rdx,[r15+358h]
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,rdi
    call mod_sub_p

mov r12,rax
    mov rcx,r12
    mov rdx,r12
    call mod_mul_p

mov r13,rax
    mov rcx,r13
    mov rdx,r12
    call mod_mul_p

mov r14,rax
    mov rcx,r12
    mov rdx,[r15+350h]
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,[r15+368h]
    call mod_mul_p

mov [r15+38h],rax
    mov rcx,rdi
    mov rdx,r13
    call mod_mul_p

mov r12,rax
    mov rcx,r12
    mov rdx,2
    call mod_mul_p

mov r13,rax
    mov rcx,r13
    mov rdx,r14
    call mod_add_p

mov r13,rax
    mov rcx,rsi
    mov rdx,[r15+360h]
    call mod_mul_p

mov rdi,rax
    mov rcx,rdi
    mov rdx,[r15+350h]
    call mod_mul_p

mov rdi,rax
    mov rcx,rdi
    mov rdx,rbp
    call mod_sub_p

mov rdi,rax
    mov rcx,rdi
    mov rdx,rdi
    call mod_mul_p

mov rsi,rax
    mov rcx,rsi
    mov rdx,r13
    call mod_sub_p

mov [r15+28h],rax
    mov rcx,r12
    mov rdx,[r15+28h]
    call mod_sub_p

mov r13,rax
    mov rcx,rdi
    mov rdx,r13
    call mod_mul_p

mov rdi,rax
    mov rcx,rbp
    mov rdx,r14
    call mod_mul_p

mov r13,rax
    mov rcx,rdi
    mov rdx,r13
    call mod_sub_p

mov [r15+30h],rax
    pop r14
    pop r13
    pop r12
    pop rbp
    pop rdi
    pop rsi
    ret
pa endp
```
点加对应的c代码(不完全对应)...
```
__declspec(noinline) void pa()
{
    DWORD64 z1_2 = mod_mul_p(ip31_z, ip31_z);
    DWORD64 z2_2 = mod_mul_p(ip32_z, ip32_z);
    DWORD64 A = mod_mul_p(ip31_x, z2_2);
    DWORD64 B = mod_mul_p(mod_mul_p(ip31_y, z2_2), ip32_z);
    DWORD64 M = mod_sub_p(mod_mul_p(ip32_x, z1_2), A);
    DWORD64 N = mod_sub_p(mod_mul_p(mod_mul_p(ip32_y, z1_2), ip31_z), B);
    DWORD64 M2 = mod_mul_p(M, M);
    DWORD64 M3 = mod_mul_p(M2, M);
    DWORD64 AM2 = mod_mul_p(A, M2);
    rp3_x = mod_sub_p(mod_mul_p(N, N), mod_add_p(M3, mod_mul_p(2, AM2)));
    rp3_y = mod_sub_p(mod_mul_p(N, mod_sub_p(AM2, rp3_x)), mod_mul_p(B, M3));
    rp3_z = mod_mul_p(mod_mul_p(M, ip31_z), ip32_z);
}
```
模n,模p 加,减,素数模乘法逆元...
```
mod_add_p proc
    mov rax,rcx
    add rax,rdx
    mov rdx,0
    mov rcx,P
    div rcx
    jnc ed
    mov rax,rdx
    add rax,PA_P
    xor rdx,rdx
    div rcx
ed:
    mov rax,rdx
    ret
mod_add_p endp

mod_add_n proc
    mov rax,rcx
    add rax,rdx
    mov rdx,0
    mov rcx,N
    div rcx
    jnc ed
    mov rax,rdx
    add rax,PA_N
    xor rdx,rdx
    div rcx
ed:
    mov rax,rdx
    ret
mod_add_n endp

mod_sub_p proc
    mov rax,rcx    
    sub rax,rdx
    mov rdx,0
    mov rcx,P
    jnc ed
    neg rax
    div rcx
    mov rax,P
    sub rax,rdx
    ret
ed:
    div rcx
    mov rax,rdx
    ret
mod_sub_p endp

mod_sub_n proc
    mov rax,rcx    
    sub rax,rdx
    mov rdx,0
    mov rcx,N
    jnc ed
    neg rax
    div rcx
    mov rax,N
    sub rax,rdx
    ret
ed:
    div rcx
    mov rax,rdx
    ret
mod_sub_n endp

pow_mod_p proc
    mov rdi,1
    mov rsi,P
    sub rsi,2
lp:
    shr rsi,1
    jnc nm
    mov rdx,rdi
    call mod_mul_p
    mov rdi,rax
nm:
    mov rdx,rcx
    call mod_mul_p
    mov rcx,rax
    test rsi,rsi
    jnz lp
    mov rax,rdi
    ret
pow_mod_p endp

pow_mod_n proc
    mov rdi,1
    mov rsi,N
    sub rsi,2
lp:
    shr rsi,1
    jnc nm
    mov rdx,rdi
    call mod_mul_n
    mov rdi,rax
nm:
    mov rdx,rcx
    call mod_mul_n
    mov rcx,rax
    test rsi,rsi
    jnz lp
    mov rax,rdi
    ret
pow_mod_n endp
```
模乘...
```
mod_mul proc;2的幂特化概率太低,去掉减少分支
    push rsi
    push rdi
    push rbp
    push r13
;    mov rbp,r15
;    dec rbp
    mulx r13,rdi,rcx
;    test rbp,r15
;    jnz ct
;    and rdi,rbp
;    mov rax,rdi
;    jmp ed
;ct:
    mov rdx,rdi
    mulx rsi,r10,r8
    mulx rbp,rax,r9
    mov rdx,r13
    mulx r8,r11,r8
    mulx r9,rax,r9
    add rbp,r10
    adc rsi,r11
    adc r8,0
    add rbp,rax
    adc rsi,r9
    adc r8,0
    mov rdx,r15
    mulx rbp,r8,r8
    mulx rsi,rbp,rsi
    add rsi,r8
    sub rdi,rbp
    sbb r13,rsi
    sub rdi,r15
    jnc ed
    add rdi,r15
    test r13,r13
    jz ed
    sub r15,rdi
    mov rdi,0ffffffffffffffffh
    sub rdi,r15
    inc rdi
ed:
    mov rax,rdi
    pop r13
    pop rbp
    pop rdi
    pop rsi
    ret
mod_mul endp

mod_mul_n proc
    push r15
    push r8
    push r9
    mov r15,N
    mov r8,N_MH
    mov r9,N_ML
    call mod_mul
    pop r9
    pop r8
    pop r15
    ret
mod_mul_n endp

mod_mul_p proc
    push r15
    push r8
    push r9
    mov r15,P
    mov r8,P_MH
    mov r9,P_ML
    call mod_mul
    pop r9
    pop r8
    pop r15
    ret
mod_mul_p endp
```
点乘标量...
```
sp_mul proc;由于分支效率问题,未对乘以0或1特化处理,所以切记不可通用!!!
;    cmp rcx,1
;    ja nbin
;    jz nz
;    mov [r15+18h],0
;    mov [r15+20h],1
;    ret
;nz:
;    mov rcx,[r15+340h]
;    mov [r15+18h],rcx
;    mov rcx,[r15+348h]
;    mov [r15+20h],rcx
;    ret
;nbin:
    push rbx
    bsr rbx,rcx
    mov rdi,rcx
    shl rdi,1
    xor rdi,rcx
    mov r8,[r15+340h]
    mov r9,[r15+348h]
    mov rsi,1
    mov [r15+350h],rsi
    call pd
    mov r10,[r15+28h]
    mov r11,[r15+30h]
    mov rax,[r15+38h]
lp:
    bt rdi,rbx
    jnc nxc
    xchg r10,r8
    xchg r11,r9
    xchg rax,rsi
nxc:
    mov [r15+340h],r10
    mov [r15+348h],r11
    mov [r15+350h],rax
    call pd
    mov [r15+358h],r8
    mov [r15+360h],r9
    mov [r15+368h],rsi
    mov r8,[r15+28h]
    mov r9,[r15+30h]
    mov rsi,[r15+38h]
    call pa
    mov r10,r8
    mov r11,r9
    mov rax,rsi
    mov r8,[r15+28h]
    mov r9,[r15+30h]
    mov rsi,[r15+38h]
    dec rbx
    jnz lp
    test rdi,1
    cmovnz r10,r8
    cmovnz r11,r9
    cmovnz rax,rsi
    mov [r15+340h],r10
    mov [r15+348h],r11
    mov [r15+350h],rax
    call point_to_affine
    pop rbx
    ret
sp_mul endp
```
kangaroo汇编版...
```
    push rcx
    push rdx
    push rsi
    push rdi
    push r8
    push r9
    push r10
    push r11

push rbx
    push rbp
    push r12
    push r13
    push r14
    push r15
    sub rsp,10h
    mov r15,rcx
    mov rbx,[r15+10h]
    mov r13,[r15+8]
    sub rbx,r13
    mov rcx,rbx
    call _sqrt
    mov rbx,rax
    inc rbx
    mov [rsp],rbx
    bsr rbx,rbx
    inc rbx
    mov dword ptr [rsp+0ch],ebx
    xor rax,rax
    mov dword ptr [rsp+8],0
lp1:
    xor r14,r14
lp2:
    mov rbx,[r15]
    mov rax,5851f42d4c957f2dh
    mul rbx
    mov rbx,rax
    inc rbx
    mov [r15],rbx
    mov r13,[rsp]
    dec r13
    xor rdx,rdx
    mov rax,rbx
    div r13
    mov rbx,rdx
    inc rbx
    cmp rbx,1
    jnz n1
    mov rbx,G_x
    lea rax,[r15+40h]
    mov [rax+r14*8],rbx
    mov rbx,G_y
    lea rax,[r15+140h]
    mov [rax+r14*8],rbx
    lea rax,[r15+240h]
    mov qword ptr [rax+r14*8],1
    jmp le1
n1:
    mov r13,G_x
    mov [r15+340h],r13
    mov r13,G_y
    mov [r15+348h],r13
    mov rcx,rbx
    call sp_mul
    mov r13,[r15+18h]
    lea rax,[r15+40h]
    mov [rax+r14*8],r13
    mov r13,[r15+20h]
    lea rax,[r15+140h]
    mov [rax+r14*8],r13
    lea rax,[r15+240h]
    mov [rax+r14*8],rbx
le1:
    inc r14
    cmp r14d,dword ptr [rsp+0ch]
    jnz lp2
    mov rbx,G_x
    mov [r15+340h],rbx
    mov rbx,G_y
    mov [r15+348h],rbx
    mov rcx,[r15+10h]
    call sp_mul
    mov rbp,[r15+20h]
    mov rax,[r15+18h]
    mov rbx,[r15+10h]
    xor r14,r14
lp3:
    xor rdx,rdx
    xor rax,rbp
    mov r13d,[rsp+0ch]
    div r13
    mov rax,rdx
    lea r13,[r15+40h]
    mov r13,[r13+rax*8]
    mov [r15+358h],r13
    lea r13,[r15+140h]
    mov r13,[r13+rax*8]
    mov [r15+360h],r13
    mov qword ptr [r15+368h],1
    lea r13,[r15+240h]
    mov r13,[r13+rax*8]
    mov [r15+348h],rbp
    mov rbp,[r15+18h]
    mov [r15+340h],rbp
    mov qword ptr [r15+350h],1
    call pa
    mov rax,[r15+28h]
    mov [r15+340h],rax
    mov rax,[r15+30h]
    mov [r15+348h],rax
    mov rax,[r15+38h]
    mov [r15+350h],rax
    call point_to_affine
    mov rax,[r15+18h]
    mov rbp,[r15+20h]
    add rbx,r13
    inc r14
    cmp r14,[rsp]
    jnz lp3
    mov r13,rax
    mov r9,rbp
    mov r12,p_y
    mov rbp,p_x
    xor r14,r14
lp4:
    mov [r15+340h],rbp
    mov [r15+348h],r12
    mov qword ptr [r15+350h],1
    mov rax,r12
    xor rdx,rdx
    xor rax,rbp
    mov r12d,[rsp+0ch]
    div r12
    mov rax,rdx
    lea r12,[r15+40h]
    mov r12,[r12+rax*8]
    mov [r15+358h],r12
    lea r12,[r15+140h]
    mov r12,[r12+rax*8]
    mov [r15+360h],r12
    mov qword ptr [r15+368h],1
    lea r12,[r15+240h]
    mov r12,[r12+rax*8]
    add r14,r12
    call pa
    mov rax,[r15+28h]
    mov [r15+340h],rax
    mov rax,[r15+30h]
    mov [r15+348h],rax
    mov rax,[r15+38h]
    mov [r15+350h],rax
    call point_to_affine
    mov rbp,[r15+18h]
    mov r12,[r15+20h]
    mov rax,rbx
    sub rax,r14
    jc lp4e
    cmp rax,[r15+8]
    jc lp4e
    cmp rax,[r15+10h]
    jnc lp4
    cmp rbp,r13
    jnz lp4
    cmp r12,r9
    jnz lp4
    jmp ed
lp4e:
    inc dword ptr [rsp+8]
    cmp dword ptr [rsp+8],8
    jnz lp1
    xor rax,rax
ed:
    add rsp,10h
    pop r15
    pop r14
    pop r13
    pop r12
    pop rbp
    pop rbx

pop r11
    pop r10
    pop r9
    pop r8
    pop rdi
    pop rsi
    pop rdx
    pop rcx
    ret
fuck_ecc_kangaroo endp
```
kangaroo c版...
```
#define TRYS 8//并行袋鼠数
#define M_SA 32//随机数池大小,和小于根号n的最大的2的幂次数有关
#define SEED_CST 0x5851f42d4c957f2d
DWORD64 M_xa[M_SA];
DWORD64 M_ya[M_SA];
DWORD64 M_za[M_SA];
DWORD64 ip31_x, ip31_y, ip31_z, ip32_x, ip32_y, ip32_z;
DWORD64 rp2_x, rp2_y, rp3_x, rp3_y, rp3_z;
typedef struct
{
    DWORD64 G_x;
    DWORD64 G_y;
    DWORD64 p_x;
    DWORD64 p_y;
    DWORD64 seed;
    DWORD64 start;
    DWORD64 end;
}PB;
__declspec(noinline) DWORD64 fuck_ecc_kangaroo(PB* pb)
{
    register DWORD64 w = pb->end - pb->start;
    DWORD64 N = _sqrt(w) + 1;
    DWORD64 m_s = 0;
    _BitScanReverse64((DWORD*)&m_s, N);//小于根号n的最大的2的幂次数
    m_s++;
    register DWORD64 i;
    register DWORD64 j;
    i = 0;
    while (1)
    {
        j = 0;
        while (1)
        {
            pb->seed = pb->seed * SEED_CST + 1;
            register DWORD64 rd = pb->seed % (N - 1) + 1;
            if (rd == 1)//特化标量1,避免点乘标量无限循环
            {
                M_xa[j] = pb->G_x;
                M_ya[j] = pb->G_y;
                M_za[j] = 1;
            }
            else
            {
                ip31_x = pb->G_x;
                ip31_y = pb->G_y;
                sp_mul(rd);
                M_xa[j] = rp2_x;
                M_ya[j] = rp2_y;
                M_za[j] = rd;
            }
            j++;
            if (j == m_s)
            {
                break;
            }
        }
        register DWORD64 H_x;
        register DWORD64 H_y;
        ip31_x = pb->G_x;
        ip31_y = pb->G_y;
        sp_mul(pb->end);
        H_x = rp2_x;
        H_y = rp2_y;
        register DWORD64 c = pb->end;
        j = 0;
        while (1)//家袋鼠跳跃
        {
            register DWORD64 ind = (H_x ^ H_y) % m_s;
            register DWORD64 xx = M_xa[ind];
            register DWORD64 yy = M_ya[ind];
            register DWORD64 zz = M_za[ind];
            ip31_x = H_x;
            ip31_y = H_y;
            ip31_z = 1;
            ip32_x = xx;
            ip32_y = yy;
            ip32_z = 1;
            pa();//不会传入无穷远点
            ip31_x = rp3_x;
            ip31_y = rp3_y;
            ip31_z = rp3_z;
            point_to_affine();
            H_x = rp2_x;
            H_y = rp2_y;
            c += zz;
            j++;
            if (j == N)
            {
                break;
            }
        }
        register DWORD64 m_x = H_x;
        register DWORD64 m_y = H_y;
        H_x = pb->p_x;
        H_y = pb->p_y;
        j = 0;
        while (1)//野袋鼠碰撞家袋鼠
        {
            register DWORD64 ind = (H_x ^ H_y) % m_s;
            register DWORD64 xx = M_xa[ind];
            register DWORD64 yy = M_ya[ind];
            register DWORD64 zz = M_za[ind];
            ip31_x = H_x;
            ip31_y = H_y;
            ip31_z = 1;
            ip32_x = xx;
            ip32_y = yy;
            ip32_z = 1;
            pa();//不会传入无穷远点
            ip31_x = rp3_x;
            ip31_y = rp3_y;
            ip31_z = rp3_z;
            point_to_affine();
            H_x = rp2_x;
            H_y = rp2_y;
            j += zz;
            register DWORD64 S = c < j;
            register DWORD64 res = c - j;
            if (S || res < pb->start)
            {
                break;
            }
            else if (res < pb->end)
            {
                if (H_x == m_x && H_y == m_y)
                {
                    return res;
                }
            }
        }
        i++;
        if (i == TRYS)//8对袋鼠并行,提高成功率
        {
            break;
        }
    }
    return 0;
}
```
由于kangaroo只有在高并发下才能发挥优势,所以上了线程池....我的机器是8核的,但是效果仍然不咋样...应该是并发数还是太少...

rho好多了,在我的机器上3秒钟能跑2^28级的数据256次,2^40的数据3次,10秒钟能跑2^48的数据一次.

c调用代码如下...
```
#include<windows.h>
#include<immintrin.h>
#define P_S 32//rho随机数池大小,越大成功率越高
#define M_S 16384//rho哈希表大小,越大结果碰撞概率越高,存储碰撞概率越小
#define TRYS 8//rho尝试次数和袋鼠对数
#define M_SA 32//kangaroo随机池大小
#define SEED_MUL 0x5851f42d4c957f2d
#define SEED_ADD 0xb5026f5aa96619e9
typedef struct
{
    DWORD64 seed;//+0h
    DWORD64 start;//+8h
    DWORD64 end;//+10h
    DWORD64 rp2_x;//+18h
    DWORD64 rp2_y;//+20h
    DWORD64 rp3_x;//+28h
    DWORD64 rp3_y;//+30h
    DWORD64 rp3_z;//+38h
    DWORD64 M_xa[M_SA];//+40h
    DWORD64 M_ya[M_SA];//+140h
    DWORD64 M_za[M_SA];//+240h
    DWORD64 ip31_x;//+340h
    DWORD64 ip31_y;//+348h
    DWORD64 ip31_z;//+350h
    DWORD64 ip32_x;//+358h
    DWORD64 ip32_y;//+360h
    DWORD64 ip32_z;//+368h
}PB;
typedef struct
{
    DWORD64 d1;
    DWORD64 d2;
    DWORD64 d3;
    DWORD64 rp2_x;
    DWORD64 rp2_y;
    DWORD64 rp3_x;
    DWORD64 rp3_y;
    DWORD64 rp3_z;  
    DWORD64 arr1[32];
    DWORD64 arr2[32];
    DWORD64 M_x[32];
    DWORD64 ip31_x;
    DWORD64 ip31_y;
    DWORD64 ip31_z;
    DWORD64 ip32_x;
    DWORD64 ip32_y;
    DWORD64 ip32_z;
    DWORD64 keys[M_S];
    DWORD64 vas[M_S];
    DWORD64 vbs[M_S];
    DWORD64 M_y[32];
}RC_CPU;
typedef int my_sprintf(char* a, size_t b, const char* c, va_list d);
extern __declspec(noinline) __m128 __fastcall _div(DWORD64 high, DWORD64 low, DWORD64 div);
extern __declspec(noinline) DWORD64 __fastcall _sqrt(DWORD64 x);
extern __declspec(noinline) DWORD64 __fastcall mod_add_n(DWORD64 a, DWORD64 b);
extern __declspec(noinline) DWORD64 __fastcall mod_sub_n(DWORD64 a, DWORD64 b);
extern __declspec(noinline) DWORD64 __fastcall mod_mul_n(DWORD64 a, DWORD64 b);
extern __declspec(noinline) DWORD64 __fastcall pow_mod_n(DWORD64 base);
extern __declspec(noinline) void __fastcall sp_mul(DWORD64 s);
extern __declspec(noinline) void point_to_affine();
extern __declspec(noinline) void pa();
extern __declspec(noinline) DWORD64 __fastcall rho_stub(DWORD64 rc, DWORD64 n);
extern __declspec(noinline) DWORD64 __fastcall fuck_ecc_kangaroo(DWORD64 pb);
extern __declspec(noinline) DWORD64 __fastcall test_n_1(DWORD64 n);
DWORD64 G_x, G_y;//基点
DWORD64 p_x, p_y;//公钥
DWORD64 A;
DWORD64 P, PA_P, P_MH, P_ML;//预计算值
DWORD64 N, PA_N, N_MH, N_ML;//预计算值
HANDLE heap = 0;
volatile LONG64 task_cnt = 0;
volatile LONG64 found_end = 0;
HANDLE done_ev = 0;
__declspec(noinline) void my_printf(const char* format, ...)
{
    DWORD i;
    char buffer[1024];
    for (i = 0; i < 1024; i++)
    {
        buffer[i] = 0;
    }
    va_list args;
    va_start(args, format);
    HMODULE ntdll = GetModuleHandleA("ntdll.dll");
    if (ntdll)
    {
        PVOID fuck_crt = GetProcAddress(ntdll, "_vsnprintf");
        if (fuck_crt)
        {
            ((my_sprintf*)fuck_crt)(buffer, sizeof(buffer), format, args);
            DWORD bytes_written;
            WriteConsoleA(GetStdHandle(STD_OUTPUT_HANDLE), buffer, lstrlenA(buffer), &bytes_written, NULL);
        }
    }
}
void print_test_time(const char* id, LARGE_INTEGER f, LARGE_INTEGER b, LARGE_INTEGER e)
{
    DWORD64 time = e.QuadPart - b.QuadPart;
    DWORD64 freq = f.QuadPart;
    register __m128 ret = _div(0, time, freq);
    DWORD64 secs = ret.m128_i64[0];
    my_printf("%s time: %llu.%llus\n", id, secs, time - secs * freq);
}
__forceinline void pre_calc(DWORD64 a, DWORD64 p, DWORD64 n)//预计算,实现无溢出模加和模乘
{
    A = a;
    P = p;
    N = n;
    PA_P = 0xffffffffffffffff % p + 1;
    PA_N = 0xffffffffffffffff % n + 1;
    register __m128 pre = _div(0xffffffffffffffff, 0xffffffffffffffff, p);
    P_ML = pre.m128_u64[0];
    P_MH = pre.m128_u64[1];
    pre = _div(0xffffffffffffffff, 0xffffffffffffffff, n);
    N_ML = pre.m128_u64[0];
    N_MH = pre.m128_u64[1];
}
__declspec(noinline) DWORD64 fuck_ecc_rho_main(RC_CPU* rc, DWORD64 n)
{
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    register DWORD64 i, j;
    DWORD tmp;
    BitScanReverse64(&tmp, n);
    DWORD64 kmsk = (1 << ((DWORD64)tmp >> 2)) - 1;//插入哈希休息间隔,为保证高效率,使用模2的幂,与运算实现
    DWORD64 reset_bound = _sqrt(n) << 3;//经过实验得出的尝试次数
    i = 0;
    while (i < TRYS)
    {
        j = 0;
        while (j < P_S)
        {
            seed = seed * SEED_MUL + SEED_ADD;
            rc->arr1[j] = seed % n;
            seed = seed * SEED_MUL + SEED_ADD;
            rc->arr2[j] = seed % n;
            rc->ip31_x = G_x;
            rc->ip31_y = G_y;
            rc->arr1[j] = rc->arr1[j] + 2;
            sp_mul(rc->arr1[j]);//传入随机数的时候+2,不让他为0或1,加上后当大于阶时候可绕回相应乘积,保证随机数均匀性,
            rc->arr1[j] = rc->arr1[j] % n;//后面由于要算加法,再模n即可,这样效率不会有大损失,这个方法只能在rho使用,因为在kangaroo要模2的幂
            register DWORD64 tmp_x = rc->rp2_x;
            register DWORD64 tmp_y = rc->rp2_y;
            rc->ip31_x = p_x;
            rc->ip31_y = p_y;
            rc->arr2[j] = rc->arr2[j] + 2;
            sp_mul(rc->arr2[j]);
            rc->arr2[j] = rc->arr2[j] % n;
            rc->ip31_x = tmp_x;
            rc->ip31_y = tmp_y;
            rc->ip31_z = 1;
            rc->ip32_x = rc->rp2_x;
            rc->ip32_y = rc->rp2_y;
            rc->ip32_z = 1;
            pa();//不会传入无穷远点
            rc->ip31_x = rc->rp3_x;
            rc->ip31_y = rc->rp3_y;
            rc->ip31_z = rc->rp3_z;
            point_to_affine();
            rc->M_x[j] = rc->rp2_x;
            rc->M_y[j] = rc->rp2_y;
            j++;
        }
        seed = seed * SEED_MUL + SEED_ADD;
        DWORD64 ax = seed % n;
        DWORD64 bx = 0;
        rc->ip31_x = G_x;
        rc->ip31_y = G_y;
        ax = ax + 2;
        sp_mul(ax);
        ax = ax % n;
        DWORD64 x_x = rc->rp2_x;
        DWORD64 x_y = rc->rp2_y;
        j = 0;
        while (j < M_S)
        {
            rc->keys[j] = 0;
            rc->vas[j] = 0;
            rc->vbs[j] = 0;
            j++;
        }
        register DWORD64 k;
        k = 0;
        while (k < reset_bound)
        {
            DWORD64 hash = x_x ^ x_y;
            DWORD64 s = hash % P_S;
            rc->ip31_x = x_x;
            rc->ip31_y = x_y;
            rc->ip31_z = 1;
            rc->ip32_x = rc->M_x[s];
            rc->ip32_y = rc->M_y[s];
            rc->ip32_z = 1;
            pa();//不会传入无穷远点
            rc->ip31_x = rc->rp3_x;
            rc->ip31_y = rc->rp3_y;
            rc->ip31_z = rc->rp3_z;
            point_to_affine();
            x_x = rc->rp2_x;
            x_y = rc->rp2_y;
            ax = mod_add_n(ax, rc->arr1[s]);
            bx = mod_add_n(bx, rc->arr2[s]);
            register DWORD64 indice = hash % M_S;
            register DWORD64 find = rc->keys[indice];
            if (find == hash)//哈希表里面存的哈希与输入值发生碰撞
            {
                DWORD64 by = rc->vbs[indice];
                if (bx != by)//避免除数为0
                {
                    DWORD64 ay = rc->vas[indice];
                    register DWORD64 res = mod_mul_n(mod_sub_n(ay, ax), pow_mod_n(mod_sub_n(bx, by)));
                    if (res > 1)//特化处理,毕竟谁把私钥设为0或1呢#(滑稽)
                    {
                        rc->ip31_x = G_x;
                        rc->ip31_y = G_y;
                        sp_mul(res);
                        if (rc->rp2_x == p_x && rc->rp2_y == p_y)
                        {
                            return res;
                        }
                    }
                }
                break;
            }
            else if ((k & kmsk) == 0)//对应位置插入哈希,如冲突则覆盖
            {              
                rc->keys[indice] = hash;
                rc->vas[indice] = ax;
                rc->vbs[indice] = bx;
            }
            k++;
        }
        i++;
    }
    return 0;
}
__declspec(noinline) DWORD64 fuck_ecc_rho(DWORD64 Gx, DWORD64 Gy, DWORD64 px, DWORD64 py, DWORD64 a, DWORD64 p, DWORD64 n)
{
    RC_CPU* rc = NULL;
    while (!(rc = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(RC_CPU))));
    G_x = Gx;
    G_y = Gy;
    p_x = px;
    p_y = py;
    pre_calc(a, p, n);
    DWORD64 ret = rho_stub((DWORD64)rc, n);
    while (!HeapFree(heap, 0, rc));
    return ret;
}
VOID CALLBACK kangaroo_worker(PTP_CALLBACK_INSTANCE Instance, PVOID Parameter, PTP_WORK Work)
{  
    UNREFERENCED_PARAMETER(Instance);
    UNREFERENCED_PARAMETER(Work);
    DWORD64 init_seed = *(DWORD64*)((DWORD64)Parameter);
    DWORD64 ret = fuck_ecc_kangaroo((DWORD64)Parameter);
    if (ret)
    {
        my_printf("found in section [%llu,%llu)! val: %llu, init_seed: %llu!\n", *(DWORD64*)((DWORD64)Parameter + 8),
            *(DWORD64*)((DWORD64)Parameter + 0x10),
            ret, init_seed);
        InterlockedExchange64(&found_end, 1);
    }
    while (!HeapFree(heap, 0, Parameter));
    if (!InterlockedDecrement64(&task_cnt))
    {
        SetEvent(done_ev);
    }
}
void fuck_ecc_kangaroo_cpu(DWORD64 Gx, DWORD64 Gy, DWORD64 px, DWORD64 py, DWORD64 a, DWORD64 p, DWORD64 n, DWORD64 pnt, DWORD64 kn)
{
    G_x = Gx;
    G_y = Gy;
    p_x = px;
    p_y = py;
    pre_calc(a, p, n);
    if (test_n_1(n))
    {
        my_printf("kangaroo: SO EASY! the val is: %llu!\n", n - 1);
        return;
    }
    my_printf("kangaroo started to run!\n");
    PTP_POOL pool = 0;
    PTP_CLEANUP_GROUP cleanupgroup = 0;
    TP_CALLBACK_ENVIRON CallBackEnviron;
    done_ev = CreateEvent(NULL, TRUE, FALSE, NULL);
    if (!done_ev)
    {
        ExitProcess(2);
    }
    InitializeThreadpoolEnvironment(&CallBackEnviron);
    pool = CreateThreadpool(NULL);
    if (!pool)
    {
        ExitProcess(3);
    }
    DWORD tn = (DWORD)(kn - 1);
    SetThreadpoolThreadMaximum(pool, tn);
    if (!SetThreadpoolThreadMinimum(pool, tn))
    {
        ExitProcess(4);
    }
    cleanupgroup = CreateThreadpoolCleanupGroup();
    if (!cleanupgroup)
    {
        ExitProcess(5);
    }
    SetThreadpoolCallbackPool(&CallBackEnviron, pool);
    SetThreadpoolCallbackCleanupGroup(&CallBackEnviron, cleanupgroup, NULL);
    pnt -= pnt % kn;//分块数对齐到线程数和2
    pnt -= pnt & 1;
    pnt += 2;
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    __m128 tmp;
    tmp = _div(0, n, pnt);
    DWORD64 psz = tmp.m128_u64[0];
    DWORD64 alignm = psz * (pnt + 1);
    tmp = _div(0, alignm, pnt);
    psz = tmp.m128_u64[0];//分块大小
    DWORD64 i, j;
    InterlockedExchange64(&found_end, 0);
    InterlockedExchange64(&task_cnt, kn);
    for (i = 0; i < pnt >> 1; i += kn)
    {
        PB* pb = NULL;
        for (j = i; j < i + kn; j++)//区间两侧同时向内执行,加快求解速度
        {
            while (!(pb = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(PB))));
            pb->seed = seed;
            pb->start = psz * j;
            pb->end = min(psz * (j + 1), n - 1);
            SubmitThreadpoolWork(CreateThreadpoolWork(kangaroo_worker, pb, &CallBackEnviron));
        }
        for (j = pnt - i - kn; j < pnt - i; j++)
        {
            while (!(pb = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(PB))));
            pb->seed = seed;
            pb->start = psz * j;
            pb->end = min(psz * (j + 1), n - 1);
            SubmitThreadpoolWork(CreateThreadpoolWork(kangaroo_worker, pb, &CallBackEnviron));
        }
        WaitForSingleObject(done_ev, INFINITE);
        CloseThreadpoolCleanupGroupMembers(cleanupgroup, FALSE, NULL);
        if (InterlockedOr64(&found_end, 0))
        {
            break;
        }
        InterlockedExchange64(&task_cnt, kn);
    }
    CloseThreadpoolCleanupGroup(cleanupgroup);
    CloseThreadpool(pool);
    CloseHandle(done_ev);
    my_printf("kangaroo stopped to run!\n");
}
void test_rho_small()
{
    LARGE_INTEGER f = { 0 };
    QueryPerformanceFrequency(&f);
    LARGE_INTEGER begin = { 0 };
    LARGE_INTEGER end = { 0 };
    QueryPerformanceCounter(&begin);
    int i;
    BOOLEAN ok = TRUE;
    for (i = 0; i < 256; i++)
    {
        DWORD64 ret = fuck_ecc_rho(191675718, 82521490,
            175473600, 2388221,
            211251998, 211252001, 211268191);
        if (ret != 162468673)
        {
            ok = FALSE;
            break;
        }
    }
    QueryPerformanceCounter(&end);
    if (ok)
    {
        print_test_time("test_rho_small", f, begin, end);
    }
    else
    {
        my_printf("test_rho_small failed!\n");
    }
}
void test_rho_medium()
{
    LARGE_INTEGER f = { 0 };
    QueryPerformanceFrequency(&f);
    LARGE_INTEGER begin = { 0 };
    LARGE_INTEGER end = { 0 };
    QueryPerformanceCounter(&begin);
    int i;
    BOOLEAN ok = TRUE;
    for (i = 0; i < 3; i++)
    {
        DWORD64 ret = fuck_ecc_rho(10213711006, 691782545293,
            550111661226, 844123301629,
            926437755434, 926437755437, 926437689103);
        if (ret != 276551929979)
        {
            ok = FALSE;
            break;
        }
    }
    QueryPerformanceCounter(&end);
    if (ok)
    {
        print_test_time("test_rho_medium", f, begin, end);
    }
    else
    {
        my_printf("test_rho_medium failed!\n");
    }
}
void test_rho_large()
{
    LARGE_INTEGER f = { 0 };
    QueryPerformanceFrequency(&f);
    LARGE_INTEGER begin = { 0 };
    LARGE_INTEGER end = { 0 };
    QueryPerformanceCounter(&begin);
    DWORD64 ret = fuck_ecc_rho(80981838096835, 11840847127648,
        71218699423971, 144038249525371,
        207699102741836, 207699102741839, 207699108502381);
    QueryPerformanceCounter(&end);
    if (ret == 87959153103470)
    {
        print_test_time("test_rho_large", f, begin, end);
    }
    else
    {
        my_printf("test_rho_large failed!\n");
    }
}
void test_idiot_kangaroo()//解为n-1时测试
{
    fuck_ecc_kangaroo_cpu(4071173633737755987, 2894341199150337603,
        4071173633737755987, 8869546445627838910,
        11763887644778176510, 11763887644778176513, 11763887647171925173,
        2, 2);//11763887647171925172
}
void test_kangaroo_sloth()
{
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    int i;
    DWORD64 a = 93127760027128;
    DWORD64 p = 93127760027131;
    DWORD64 n = 93127766672107;
    G_x = 78289974459711;
    G_y = 63568046851464;
    p_x = 38019674211203;
    p_y = 35079481704560;
    pre_calc(a, p, n);
    DWORD64 pnt = 33333363;
    __m128 tmp;
    tmp = _div(0, n, pnt);
    DWORD64 psz = tmp.m128_u64[0];
    DWORD64 alignm = psz * (pnt + 1);
    tmp = _div(0, alignm, pnt);
    psz = tmp.m128_u64[0];
    for (i = 11162854 - 333; i < 11162854 + 222; i++)
    {
        PB* pb = NULL;
        while (!(pb = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(PB))));
        pb->seed = seed;
        pb->start = psz * i;
        pb->end = min(psz * (i + 1), n - 1);
        DWORD64 ret = fuck_ecc_kangaroo((DWORD64)pb);
        while (!HeapFree(heap, 0, pb));
        if (ret)
        {
            my_printf("found at section %d!val:%llu!", i, ret);
            if (ret == 31187151843342)
            {
                my_printf("correct!\n");
            }
            else
            {
                my_printf("wrong!\n");
            }
        }
    }
    my_printf("ended!\n");
}
void test_kangaroo_cpu_small()
{
    fuck_ecc_kangaroo_cpu(191675718, 82521490,
        175473600, 2388221,
        211251998, 211252001, 211268191,
        256, 128);
}
void test_kangaroo_cpu_medium()
{
    fuck_ecc_kangaroo_cpu(10213711006, 691782545293,
        550111661226, 844123301629,
        926437755434, 926437755437, 926437689103,
        256, 128);
}
void test_kangaroo_cpu_large()
{
    fuck_ecc_kangaroo_cpu(80981838096835, 11840847127648,
        71218699423971, 144038249525371,
        207699102741836, 207699102741839, 207699108502381,
        256, 128);
}
void test_cpu_rho()
{
    test_rho_small();
    test_rho_medium();
    test_rho_large();
}
void test_cpu_kangaroo()
{
    test_idiot_kangaroo();
    test_kangaroo_sloth();
    test_kangaroo_cpu_small();
    test_kangaroo_cpu_medium();
    test_kangaroo_cpu_large();
}
void test_cpu()
{
    test_cpu_rho();
    test_cpu_kangaroo();
}
__declspec(noreturn) void zxmain()
{
    heap = GetProcessHeap();
    test_cpu();
    my_printf("done!!!\r\n");
    FatalExit(0);
}
```
然而,线程池kangaroo的效率和单线程rho还是差得远.所以考虑上gpu,进一步提升并发数.

关于异构计算,由于我刚入门不久,以下内容可能存在大量理论错误,请巨佬们轻喷!

我的gpu是nvidia tesla p4云计算卡,cuda版本12.0,显存8gb,为pascal架构,最大支持32*32个线程(由于warp大小为32,所以x,y应均为32倍数,避免多余线程,又由于线程块大小在每个方向均有限制,所以分量大小应相等以取得最大值).

先将kangaroo改为ptx版本,观察到效率提升10倍左右,但是与cpu单线程的rho还有一段差距.我认识到,在单卡这样的并发规模下,kangaroo仍然不能得到最好的发挥.

所以我放弃优化kangaroo,转而将rho改为适合并行的版本,先给出c版单线程迭代代码(每次循环代表一个线程,16次尝试,每次尝试迭代256个分区,共16*256=4096个任务,每个线程执行4个任务),再由该代码转为ptx代码.
```
__declspec(noinline) DWORD64 fuck_ecc_rho_main_1(RC_CPU* rc, DWORD64 n)//n>=2^8
{
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    register DWORD64 i, j;
    DWORD tmp;
    BitScanReverse64(&tmp, n);
    DWORD64 kmsk = (1 << ((DWORD64)tmp >> 2)) - 1;
    DWORD64 sqn = _sqrt(n);
    DWORD64 bnd = (sqn >> 4) + (sqn >> 6);//实验得出的bounding
    i = 0;
    while (i < 16)//增加到16次尝试
    {
        j = 0;
        while (j < 32)
        {
            seed = seed * SEED_MUL + SEED_ADD;
            rc->arr1[j] = seed % n;
            seed = seed * SEED_MUL + SEED_ADD;
            rc->arr2[j] = seed % n;
            rc->ip31_x = G_x;
            rc->ip31_y = G_y;
            rc->arr1[j] = rc->arr1[j] + 2;
            sp_mul(rc->arr1[j]);
            rc->arr1[j] = rc->arr1[j] % n;
            register DWORD64 tmp_x = rc->rp2_x;
            register DWORD64 tmp_y = rc->rp2_y;
            rc->ip31_x = p_x;
            rc->ip31_y = p_y;
            rc->arr2[j] = rc->arr2[j] + 2;
            sp_mul(rc->arr2[j]);
            rc->arr2[j] = rc->arr2[j] % n;
            rc->ip31_x = tmp_x;
            rc->ip31_y = tmp_y;
            rc->ip31_z = 1;
            rc->ip32_x = rc->rp2_x;
            rc->ip32_y = rc->rp2_y;
            rc->ip32_z = 1;
            pa();
            rc->ip31_x = rc->rp3_x;
            rc->ip31_y = rc->rp3_y;
            rc->ip31_z = rc->rp3_z;
            point_to_affine();
            rc->M_x[j] = rc->rp2_x;
            rc->M_y[j] = rc->rp2_y;
            j++;
        }
        register DWORD64 m;
        m = 0;
        while (m < 256)//256个区间分块
        {
            seed = seed * SEED_MUL + SEED_ADD;
            DWORD64 ax = seed % n;
            DWORD64 bx = 0;
            rc->ip31_x = G_x;
            rc->ip31_y = G_y;
            ax = ax + 2;
            sp_mul(ax);
            ax = ax % n;
            DWORD64 x_x = rc->rp2_x;
            DWORD64 x_y = rc->rp2_y;
            j = 0;
            while (j < M_S)
            {
                rc->keys[j] = 0;
                rc->vas[j] = 0;
                rc->vbs[j] = 0;
                j++;
            }
            register DWORD64 k;
            k = 0;
            while (k < bnd)
            {
                DWORD64 hash = x_x ^ x_y;
                DWORD64 s = hash % 32;
                rc->ip31_x = x_x;
                rc->ip31_y = x_y;
                rc->ip31_z = 1;
                rc->ip32_x = rc->M_x[s];
                rc->ip32_y = rc->M_y[s];
                rc->ip32_z = 1;
                pa();
                rc->ip31_x = rc->rp3_x;
                rc->ip31_y = rc->rp3_y;
                rc->ip31_z = rc->rp3_z;
                point_to_affine();
                x_x = rc->rp2_x;
                x_y = rc->rp2_y;
                ax = mod_add_n(ax, rc->arr1[s]);
                bx = mod_add_n(bx, rc->arr2[s]);
                register DWORD64 indice = hash % M_S;
                register DWORD64 find = rc->keys[indice];
                if (find == hash)
                {
                    DWORD64 by = rc->vbs[indice];
                    if (bx != by)
                    {
                        DWORD64 ay = rc->vas[indice];
                        register DWORD64 res = mod_mul_n(mod_sub_n(ay, ax), pow_mod_n(mod_sub_n(bx, by)));
                        if (res > 1)//特化处理,毕竟谁把私钥设为0或1呢#(滑稽)
                        {
                            rc->ip31_x = G_x;
                            rc->ip31_y = G_y;
                            sp_mul(res);
                            if (rc->rp2_x == p_x && rc->rp2_y == p_y)
                            {
                                return res;
                            }
                        }
                    }
                    break;
                }
                else if ((k & kmsk) == 0)
                {
                    rc->keys[indice] = hash;
                    rc->vas[indice] = ax;
                    rc->vbs[indice] = bx;
                }
                k++;
            }
            m++;
        }
        i++;
    }
    return 0;
}
```
cpu端代码.
```
#include<windows.h>
#include<immintrin.h>
#include<cuda.h>
#include<cuda_runtime.h>
#define P_S 32
#define M_S 16384
#define TRYS 8
#define M_SA 32
#define SEED_MUL 0x5851f42d4c957f2d
#define SEED_ADD 0xb5026f5aa96619e9
#define CUDA_CALL(x) {CUresult res;if((res=(x))!=CUDA_SUCCESS){ExitProcess(res);}}
typedef int my_sprintf(char* a, size_t b, const char* c, va_list d);
extern __declspec(noinline) __m128 __fastcall _div(DWORD64 high, DWORD64 low, DWORD64 div);
extern __declspec(noinline) DWORD64 __fastcall _sqrt(DWORD64 x);
DWORD64 G_x, G_y;
DWORD64 p_x, p_y;
DWORD64 A;
DWORD64 P, PA_P, P_MH, P_ML;
DWORD64 N, PA_N, N_MH, N_ML;
HANDLE heap = 0;
HANDLE done_ev = 0;
__declspec(noinline) void my_printf(const char* format, ...)
{
    DWORD i;
    char buffer[1024];
    for (i = 0; i < 1024; i++)
    {
        buffer[i] = 0;
    }
    va_list args;
    va_start(args, format);
    HMODULE ntdll = GetModuleHandleA("ntdll.dll");
    if (ntdll)
    {
        PVOID fuck_crt = GetProcAddress(ntdll, "_vsnprintf");
        if (fuck_crt)
        {
            ((my_sprintf*)fuck_crt)(buffer, sizeof(buffer), format, args);
            DWORD bytes_written;
            WriteConsoleA(GetStdHandle(STD_OUTPUT_HANDLE), buffer, lstrlenA(buffer), &bytes_written, NULL);
        }
    }
}
void result_monitor(DWORD64* args)
{
    CUcontext* context = (CUcontext*)args[0];
    DWORD64* ret_uc = (DWORD64*)args[1];
    const char* id = (const char*)args[2];
    BOOLEAN printed = FALSE;
    *ret_uc = 0;
    CUDA_CALL(cuCtxSetCurrent(*context));
    while (1)
    {
        Sleep(3);
        DWORD64 val = *ret_uc;
        if (val)
        {
            if (val != 0xffffffffffffffff)
            {
                if (!printed)
                {
                    my_printf("gpu_%s: the val is %llu!\n", id, val);
                    printed = TRUE;
                }
            }
            else
            {
                break;
            }
        }
    }
    CUDA_CALL(cuMemFreeHost(ret_uc));
    CUDA_CALL(cuCtxDestroy(*context));
    while (!HeapFree(heap, 0, context));
}
void gpu_rho(char* ptx_code, DWORD64 gx, DWORD64 gy, DWORD64 px, DWORD64 py, DWORD64 aa, DWORD64 pp, DWORD64 nn)
{
    G_x = gx;
    G_y = gy;
    p_x = px;
    p_y = py;
    pre_calc(aa, pp, nn);
    if (test_n_1(nn))
    {
        my_printf("rho: SO EASY! the val is: %llu!\n", nn - 1);
        return;
    }
    DWORD64 sqn = _sqrt(nn);
    DWORD64 T = (sqn >> 4) + (sqn >> 7);
    CUdevice device;
    CUDA_CALL(cuDeviceGet(&device, 0));
    CUcontext* context;
    while (!(context = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(CUcontext))));
    CUDA_CALL(cuCtxCreate(context, 0, device));
    CUDA_CALL(cuCtxSetCurrent(*context));
    CUmodule module;
    CUDA_CALL(cuModuleLoadData(&module, ptx_code));
    CUfunction kernel;
    DWORD64 tx = 16, ty = 16;
    DWORD64 tn = tx * ty;
    DWORD64 gn1 = (DWORD64)(512 / tn), gn2 = (DWORD64)(4096 / tn);
    size_t tmpsz;
    CUdeviceptr ret_u, mx, my, keys, vas, vbs;
    volatile DWORD64* ret_u_host = 0;
    DWORD64 buf_sz = 16 * 32;
    DWORD64 buf_sz2 = 16 * 256 * M_S;
    DWORD bsr;
    BitScanReverse64(&bsr, nn);
    DWORD64 kmsk = (1 << ((DWORD64)bsr >> 2)) - 1;
    void* args[] = { &ret_u };
    CUDA_CALL(cuMemAlloc(&mx, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemAlloc(&my, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemAlloc(&keys, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemAlloc(&vas, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemAlloc(&vbs, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemsetD8(keys, 0, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemsetD8(vas, 0, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemsetD8(vbs, 0, buf_sz2 * sizeof(DWORD64)));
    CUDA_CALL(cuMemHostAlloc((DWORD64**)&ret_u_host, sizeof(DWORD64), CU_MEMHOSTALLOC_DEVICEMAP));
    CUDA_CALL(cuMemHostGetDevicePointer(&ret_u, (DWORD64*)ret_u_host, 0));
    CUdeviceptr AA, PP, NN, PAPAP, MHMHP, MLMLP, PAPAN, MHMHN, MLMLN, TT,
        gxgx, gygy, pxpx, pypy, keyskeys, vasvas, vbsvbs, ar1, ar2, mxmx, mymy, msk;
    CUDA_CALL(cuModuleGetGlobal(&AA, &tmpsz, module, "A"));
    CUDA_CALL(cuModuleGetGlobal(&PP, &tmpsz, module, "P"));
    CUDA_CALL(cuModuleGetGlobal(&NN, &tmpsz, module, "N"));
    CUDA_CALL(cuModuleGetGlobal(&PAPAP, &tmpsz, module, "PA_P"));
    CUDA_CALL(cuModuleGetGlobal(&MHMHP, &tmpsz, module, "MH_P"));
    CUDA_CALL(cuModuleGetGlobal(&MLMLP, &tmpsz, module, "ML_P"));
    CUDA_CALL(cuModuleGetGlobal(&PAPAN, &tmpsz, module, "PA_N"));
    CUDA_CALL(cuModuleGetGlobal(&MHMHN, &tmpsz, module, "MH_N"));
    CUDA_CALL(cuModuleGetGlobal(&MLMLN, &tmpsz, module, "ML_N"));
    CUDA_CALL(cuModuleGetGlobal(&TT, &tmpsz, module, "T"));
    CUDA_CALL(cuModuleGetGlobal(&gxgx, &tmpsz, module, "G_x"));
    CUDA_CALL(cuModuleGetGlobal(&gygy, &tmpsz, module, "G_y"));
    CUDA_CALL(cuModuleGetGlobal(&pxpx, &tmpsz, module, "p_x"));
    CUDA_CALL(cuModuleGetGlobal(&pypy, &tmpsz, module, "p_y"));
    CUDA_CALL(cuModuleGetGlobal(&keyskeys, &tmpsz, module, "keys"));
    CUDA_CALL(cuModuleGetGlobal(&vasvas, &tmpsz, module, "vas"));
    CUDA_CALL(cuModuleGetGlobal(&vbsvbs, &tmpsz, module, "vbs"));
    CUDA_CALL(cuModuleGetGlobal(&ar1, &tmpsz, module, "arr1"));
    CUDA_CALL(cuModuleGetGlobal(&ar2, &tmpsz, module, "arr2"));
    CUDA_CALL(cuModuleGetGlobal(&mxmx, &tmpsz, module, "m_x"));
    CUDA_CALL(cuModuleGetGlobal(&mymy, &tmpsz, module, "m_y"));
    CUDA_CALL(cuModuleGetGlobal(&msk, &tmpsz, module, "MSK"));
    CUDA_CALL(cuMemcpyHtoD(AA, &aa, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(PP, &pp, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(NN, &nn, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(PAPAP, &PA_P, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MHMHP, &P_MH, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MLMLP, &P_ML, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(PAPAN, &PA_N, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MHMHN, &N_MH, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MLMLN, &N_ML, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(TT, &T, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(gxgx, &G_x, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(gygy, &G_y, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(pxpx, &p_x, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(pypy, &p_y, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(keyskeys, &keys, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(vasvas, &vas, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(vbsvbs, &vbs, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(mxmx, &mx, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(mymy, &my, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(msk, &kmsk, sizeof(DWORD64)));
    DWORD tid;
    void* args_thr[] = { context, (DWORD64*)ret_u_host, "rho" };
    HANDLE thr = CreateThread(NULL, 0, (LPTHREAD_START_ROUTINE)result_monitor, args_thr, 0, &tid);//结果监视器
    if (thr)
    {
        CloseHandle(thr);
        thr = NULL;
    }
    DWORD64 i;
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    DWORD64 *arr1_ptr = 0, *arr2_ptr = 0, *mx_ptr = 0, *my_ptr = 0;
    while (!(arr1_ptr = HeapAlloc(heap, HEAP_ZERO_MEMORY, buf_sz * sizeof(DWORD64))));
    while (!(arr2_ptr = HeapAlloc(heap, HEAP_ZERO_MEMORY, buf_sz * sizeof(DWORD64))));
    while (!(mx_ptr = HeapAlloc(heap, HEAP_ZERO_MEMORY, buf_sz * sizeof(DWORD64))));
    while (!(my_ptr = HeapAlloc(heap, HEAP_ZERO_MEMORY, buf_sz * sizeof(DWORD64))));
    register DWORD64 tmp;
    for (i = 0; i < buf_sz; i++)//预计算随机数
    {
        seed = seed * SEED_MUL + SEED_ADD;
        tmp = (seed % nn) + 2;
        mx_ptr[i] = tmp;
        tmp %= nn;
        arr1_ptr[i] = tmp;
        seed = seed * SEED_MUL + SEED_ADD;
        tmp = (seed % nn) + 2;
        my_ptr[i] = tmp;
        tmp %= nn;
        arr2_ptr[i] = tmp;
    }
    CUDA_CALL(cuMemcpyHtoD(ar1, arr1_ptr, buf_sz * sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(ar2, arr2_ptr, buf_sz * sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(mx, mx_ptr, buf_sz * sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(my, my_ptr, buf_sz * sizeof(DWORD64)));
    while (!HeapFree(heap, 0, my_ptr));
    while (!HeapFree(heap, 0, mx_ptr));
    while (!HeapFree(heap, 0, arr2_ptr));
    while (!HeapFree(heap, 0, arr1_ptr));
    CUDA_CALL(cuModuleGetFunction(&kernel, module, "init_rho"));
    CUDA_CALL(cuLaunchKernel(kernel, (DWORD)gn1, 1, 1, (DWORD)tn, 1, 1, 0, NULL, NULL, NULL));
    CUDA_CALL(cuCtxSynchronize());
    CUDA_CALL(cuModuleGetFunction(&kernel, module, "rho"));
    CUDA_CALL(cuLaunchKernel(kernel, (DWORD)gn2, 1, 1, (DWORD)tn, 1, 1, 0, NULL, args, NULL));
    CUDA_CALL(cuCtxSynchronize());
    CUDA_CALL(cuMemFree(vbs));
    CUDA_CALL(cuMemFree(vas));
    CUDA_CALL(cuMemFree(keys));
    CUDA_CALL(cuMemFree(my));
    CUDA_CALL(cuMemFree(mx));
    CUDA_CALL(cuModuleUnload(module));
    InterlockedExchange64(ret_u_host, 0xffffffffffffffff);
}
void gpu_kangaroo(char* ptx_code, DWORD64 gx, DWORD64 gy, DWORD64 px, DWORD64 py, DWORD64 aa, DWORD64 pp, DWORD64 nn)
{
    G_x = gx;
    G_y = gy;
    p_x = px;
    p_y = py;
    pre_calc(aa, pp, nn);
    if (test_n_1(nn))
    {
        my_printf("kangaroo: SO EASY! the val is: %llu!\n", nn - 1);
        return;
    }
    CUdevice device;
    CUDA_CALL(cuDeviceGet(&device, 0));
    CUcontext* context;
    while (!(context = HeapAlloc(heap, HEAP_ZERO_MEMORY, sizeof(CUcontext))));
    CUDA_CALL(cuCtxCreate(context, 0, device));
    CUDA_CALL(cuCtxSetCurrent(*context));
    CUmodule module;
    CUDA_CALL(cuModuleLoadData(&module, ptx_code));
    CUfunction kernel;
    DWORD64 bna = 23;//bn参数
    DWORD64 tx = 16, ty = 16;
    DWORD64 tn = tx * ty;
    DWORD64 bn = bna * tn;
    DWORD64 t_bn = TRYS * bn;
    DWORD64 gn = TRYS * bna;
    __m128 tmp;
    tmp = _div(0, nn, bn);
    DWORD64 psz = tmp.m128_u64[0];
    DWORD64 alignm = psz * (bn + 1);
    tmp = _div(0, alignm, bn);
    psz = tmp.m128_u64[0];//限制:(256<=bn<=0xffffffff),n必须小于0xff00000000000000
    DWORD64 T = _sqrt(psz) + 1;
    DWORD TMP;
    BitScanReverse64(&TMP, T);
    DWORD64 K = (DWORD64)TMP;
    DWORD64 m_sz = 3 * K * t_bn;
    size_t tmpsz;
    CUdeviceptr ret_u, M;
    volatile DWORD64* ret_u_host = 0;
    void* args[] = { &ret_u };
    CUDA_CALL(cuMemAlloc(&M, m_sz * sizeof(DWORD64)));
    CUDA_CALL(cuMemHostAlloc((DWORD64**)&ret_u_host, sizeof(DWORD64), CU_MEMHOSTALLOC_DEVICEMAP));
    CUDA_CALL(cuMemHostGetDevicePointer(&ret_u, (DWORD64*)ret_u_host, 0));
    CUdeviceptr AA, PP, NN, PAPAP, MHMHP, MLMLP, gxgx, gygy, pxpx, pypy, KK, TT, MM, QQ;
    CUDA_CALL(cuModuleGetGlobal(&AA, &tmpsz, module, "A"));
    CUDA_CALL(cuModuleGetGlobal(&PP, &tmpsz, module, "P"));
    CUDA_CALL(cuModuleGetGlobal(&NN, &tmpsz, module, "N"));
    CUDA_CALL(cuModuleGetGlobal(&PAPAP, &tmpsz, module, "PA_P"));
    CUDA_CALL(cuModuleGetGlobal(&MHMHP, &tmpsz, module, "MH_P"));
    CUDA_CALL(cuModuleGetGlobal(&MLMLP, &tmpsz, module, "ML_P"));
    CUDA_CALL(cuModuleGetGlobal(&gxgx, &tmpsz, module, "G_x"));
    CUDA_CALL(cuModuleGetGlobal(&gygy, &tmpsz, module, "G_y"));
    CUDA_CALL(cuModuleGetGlobal(&pxpx, &tmpsz, module, "p_x"));
    CUDA_CALL(cuModuleGetGlobal(&pypy, &tmpsz, module, "p_y"));
    CUDA_CALL(cuModuleGetGlobal(&KK, &tmpsz, module, "K"));
    CUDA_CALL(cuModuleGetGlobal(&TT, &tmpsz, module, "T"));
    CUDA_CALL(cuModuleGetGlobal(&MM, &tmpsz, module, "M"));
    CUDA_CALL(cuModuleGetGlobal(&QQ, &tmpsz, module, "Q"));
    CUDA_CALL(cuMemcpyHtoD(AA, &aa, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(PP, &pp, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(NN, &nn, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(PAPAP, &PA_P, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MHMHP, &P_MH, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MLMLP, &P_ML, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(gxgx, &G_x, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(gygy, &G_y, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(pxpx, &p_x, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(pypy, &p_y, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(KK, &K, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(TT, &T, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(MM, &M, sizeof(DWORD64)));
    CUDA_CALL(cuMemcpyHtoD(QQ, &psz, sizeof(DWORD64)));
    DWORD tid;
    void* args_thr[] = { context, (DWORD64*)ret_u_host, "kangaroo" };
    HANDLE thr = CreateThread(NULL, 0, (LPTHREAD_START_ROUTINE)result_monitor, args_thr, 0, &tid);
    if (thr)
    {
        CloseHandle(thr);
        thr = NULL;
    }
    DWORD64 i;
    DWORD64 seed = SEED_ADD;
    seed *= SEED_MUL;
    seed *= __rdtsc();
    DWORD64* M_host = 0;
    while (!(M_host = HeapAlloc(heap, HEAP_ZERO_MEMORY, m_sz * sizeof(DWORD64))));
    for (i = 0; i < m_sz; i += 3)//预计算随机数,填入一段未初始化的全局内存,后面使用并覆盖
    {
        seed = seed * SEED_MUL + SEED_ADD;
        M_host[i] = seed;
    }
    CUDA_CALL(cuMemcpyHtoD(M, M_host, m_sz * sizeof(DWORD64)));
    while (!HeapFree(heap, 0, M_host));
    CUDA_CALL(cuModuleGetFunction(&kernel, module, "init_M"));
    CUDA_CALL(cuLaunchKernel(kernel, (DWORD)(gn * K), 1, 1, (DWORD)tn, 1, 1, 0, NULL, NULL, NULL));
    CUDA_CALL(cuCtxSynchronize());
    CUDA_CALL(cuModuleGetFunction(&kernel, module, "kangaroo"));
    CUDA_CALL(cuLaunchKernel(kernel, (DWORD)gn, 1, 1, (DWORD)tn, 1, 1, 0, NULL, args, NULL));
    CUDA_CALL(cuCtxSynchronize());
    CUDA_CALL(cuMemFree(M));
    CUDA_CALL(cuModuleUnload(module));
    InterlockedExchange64(ret_u_host, 0xffffffffffffffff);
}
void test_gpu()
{
    CUDA_CALL(cuInit(0));
    HANDLE hFile = INVALID_HANDLE_VALUE;
    DWORD bytes_read = 0;
    char* ptx_code = NULL;
    hFile = CreateFileA("C:\\Users\\n00bzx\\source\\repos\\chengkende\\kangaroo_and_rho.ptx",
        GENERIC_READ, FILE_SHARE_READ, NULL, OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, NULL);
    if (hFile == INVALID_HANDLE_VALUE)
    {
        ExitProcess(6);
    }
    LARGE_INTEGER file_size;
    if (!GetFileSizeEx(hFile, &file_size))
    {
        ExitProcess(7);
    }
    DWORD64 fs = file_size.QuadPart;
    while (!(ptx_code = HeapAlloc(heap, HEAP_ZERO_MEMORY, fs + 1)));
    if (!ReadFile(hFile, ptx_code, (DWORD)fs, &bytes_read, NULL))
    {
        ExitProcess(8);
    }
    ptx_code[fs] = 0;
    CloseHandle(hFile);
    DWORD64 i = 0;
    for (i = 0; i < 3; i++)
    {
        gpu_rho(ptx_code, 272035190, 295528426, 373098171, 362966399, 1133936968, 1133936971, 1133916691);//220476114
        gpu_kangaroo(ptx_code, 272035190, 295528426, 373098171, 362966399, 1133936968, 1133936971, 1133916691);//220476114
    }
    gpu_rho(ptx_code, 10213711006, 691782545293, 550111661226, 844123301629, 926437755434, 926437755437, 926437689103);//276551929979
    gpu_kangaroo(ptx_code, 10213711006, 691782545293, 550111661226, 844123301629, 926437755434, 926437755437, 926437689103);//276551929979
    while (!HeapFree(heap, 0, ptx_code));
}
__declspec(noreturn) void zxmain()
{
    heap = GetProcessHeap();
    test_gpu();
    my_printf("done!!!\r\n");
    FatalExit(0);
}
```
gpu端ptx代码...
```
.version 6.0
.target sm_61
.const .u64 A, P, N, PA_P, MH_P, ML_P, PA_N, MH_N, ML_N, G_x, G_y, p_x, p_y, K, T, M, Q, MSK;
.const .u64 keys, vas, vbs;
.const .u64 arr1[512], arr2[512];
.const .u64 m_x, m_y;
.func (.reg .u64 out) mod_add_p (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r1, %r2, %r3;
    .reg .pred %c;
    mov.u64 %r1, a;
    mov.u64 %r2, b;
    add.cc.u64 %r1, %r1, %r2;
    addc.cc.u64 %r3, 0, 0;
    setp.eq.u64 %c, %r3, 0;
    ld.const.u64 %r3, [P];
    rem.u64 %r1, %r1, %r3;
    @%c bra ed;
    ld.const.u64 %r2, [PA_P];
    add.u64 %r1, %r1, %r2;
    rem.u64 %r1, %r1, %r3;
ed:
    mov.u64 out, %r1;
    ret;
}
.func (.reg .u64 out) mod_sub_p (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r1, %r2, %r3;
    .reg .pred %c;
    mov.u64 %r1, a;
    mov.u64 %r2, b;
    ld.const.u64 %r3, [P];
    setp.ge.u64 %c, %r1, %r2;
    sub.u64 %r1, %r1, %r2;
    @%c bra ed;
    not.b64 %r1, %r1;
    add.u64 %r1, %r1, 1;
    rem.u64 %r1, %r1, %r3;
    sub.u64 %r1, %r3, %r1;
    mov.u64 out, %r1;
    ret;
ed:
    rem.u64 %r1, %r1, %r3;
    mov.u64 out, %r1;
    ret;
}
.func (.reg .u64 out) mod_mul_p (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9, %r10, %r11;
    .reg .pred %c1, %c2;
    mov.u64 %r0, a;
    mov.u64 %r1, b;
    ld.const.u64 %r2, [P];
    ld.const.u64 %r3, [MH_P];
    ld.const.u64 %r4, [ML_P];
    sub.u64 %r5, %r2, 1;
    mad.lo.u64 %r6, %r0, %r1, 0;
    mad.hi.u64 %r7, %r0, %r1, 0;
    mad.lo.u64 %r10, %r3, %r6, 0;
    mad.hi.u64 %r8, %r3, %r6, 0;
    mad.lo.u64 %r9, %r4, %r6, 0;
    mad.hi.u64 %r5, %r4, %r6, 0;
    mad.lo.u64 %r11, %r3, %r7, 0;
    mad.hi.u64 %r3, %r3, %r7, 0;
    mad.lo.u64 %r9, %r4, %r7, 0;
    mad.hi.u64 %r4, %r4, %r7, 0;
    add.cc.u64 %r5, %r5, %r10;
    addc.cc.u64 %r8, %r8, %r11;
    addc.cc.u64 %r3, %r3, 0;
    add.cc.u64 %r5, %r5, %r9;
    addc.cc.u64 %r8, %r8, %r4;
    addc.cc.u64 %r3, %r3, 0;
    mad.lo.u64 %r5, %r8, %r2, 0;
    mad.hi.u64 %r8, %r8, %r2, 0;
    mad.lo.u64 %r3, %r3, %r2, 0;
    add.u64 %r8, %r8, %r3;
    sub.cc.u64 %r6, %r6, %r5;
    subc.cc.u64 %r7, %r7, %r8;
    setp.eq.u64 %c2, %r7, 1;
    setp.lt.u64 %c1, %r6, %r2;
    sub.u64 %r6, %r6, %r2;
    mov.u64 %r1, %r6;
    add.u64 %r6, %r6, %r2;
    mov.u64 %r3, %r6;
    sub.u64 %r2, %r2, %r6;
    sub.u64 %r4, 0xffffffffffffffff, %r2;
    add.u64 %r4, %r4, 1;
    selp.u64 %r5, %r3, %r1, %c1;
    selp.u64 %r5, %r4, %r5, %c2;
    mov.u64 out, %r5;
    ret;
}
.func (.reg .u64 out) fpm_p (.reg .u64 base)
{
    .reg .u64 %r0, %r1, %r2;
    .reg .pred %c;
    mov.u64 %r0, base;
    ld.const.u64 %r2, [P];
    sub.u64 %r2, %r2, 2;
    mov.u64 out, 1;
lp:
    and.b64 %r1, %r2, 1;
    setp.eq.u64 %c, %r1, 0;
    @%c bra nm;
    call (out), mod_mul_p, (%r0, out);
nm:
    call (%r0), mod_mul_p, (%r0, %r0);
    setp.ne.u64 %c, %r2, 0;
    shr.u64 %r2, %r2, 1;
    @%c bra lp;
    ret;
}
.func (.reg .u64 out) mod_add_n (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r1, %r2, %r3;
    .reg .pred %c;
    mov.u64 %r1, a;
    mov.u64 %r2, b;
    add.cc.u64 %r1, %r1, %r2;
    addc.cc.u64 %r3, 0, 0;
    setp.eq.u64 %c, %r3, 0;
    ld.const.u64 %r3, [N];
    rem.u64 %r1, %r1, %r3;
    @%c bra ed;
    ld.const.u64 %r2, [PA_N];
    add.u64 %r1, %r1, %r2;
    rem.u64 %r1, %r1, %r3;
ed:
    mov.u64 out, %r1;
    ret;
}
.func (.reg .u64 out) mod_sub_n (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r1, %r2, %r3;
    .reg .pred %c;
    mov.u64 %r1, a;
    mov.u64 %r2, b;
    ld.const.u64 %r3, [N];
    setp.ge.u64 %c, %r1, %r2;
    sub.u64 %r1, %r1, %r2;
    @%c bra ed;
    not.b64 %r1, %r1;
    add.u64 %r1, %r1, 1;
    rem.u64 %r1, %r1, %r3;
    sub.u64 %r1, %r3, %r1;
    mov.u64 out, %r1;
    ret;
ed:
    rem.u64 %r1, %r1, %r3;
    mov.u64 out, %r1;
    ret;
}
.func (.reg .u64 out) mod_mul_n (.reg .u64 a, .reg .u64 b)
{
    .reg.u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9, %r10, %r11;
    .reg .pred %c1, %c2;
    mov.u64 %r0, a;
    mov.u64 %r1, b;
    ld.const.u64 %r2, [N];
    ld.const.u64 %r3, [MH_N];
    ld.const.u64 %r4, [ML_N];
    sub.u64 %r5, %r2, 1;
    mad.lo.u64 %r6, %r0, %r1, 0;
    mad.hi.u64 %r7, %r0, %r1, 0;
    mad.lo.u64 %r10, %r3, %r6, 0;
    mad.hi.u64 %r8, %r3, %r6, 0;
    mad.lo.u64 %r9, %r4, %r6, 0;
    mad.hi.u64 %r5, %r4, %r6, 0;
    mad.lo.u64 %r11, %r3, %r7, 0;
    mad.hi.u64 %r3, %r3, %r7, 0;
    mad.lo.u64 %r9, %r4, %r7, 0;
    mad.hi.u64 %r4, %r4, %r7, 0;
    add.cc.u64 %r5, %r5, %r10;
    addc.cc.u64 %r8, %r8, %r11;
    addc.cc.u64 %r3, %r3, 0;
    add.cc.u64 %r5, %r5, %r9;
    addc.cc.u64 %r8, %r8, %r4;
    addc.cc.u64 %r3, %r3, 0;
    mad.lo.u64 %r5, %r8, %r2, 0;
    mad.hi.u64 %r8, %r8, %r2, 0;
    mad.lo.u64 %r3, %r3, %r2, 0;
    add.u64 %r8, %r8, %r3;
    sub.cc.u64 %r6, %r6, %r5;
    subc.cc.u64 %r7, %r7, %r8;
    setp.eq.u64 %c2, %r7, 1;
    setp.lt.u64 %c1, %r6, %r2;
    sub.u64 %r6, %r6, %r2;
    mov.u64 %r1, %r6;
    add.u64 %r6, %r6, %r2;
    mov.u64 %r3, %r6;
    sub.u64 %r2, %r2, %r6;
    sub.u64 %r4, 0xffffffffffffffff, %r2;
    add.u64 %r4, %r4, 1;
    selp.u64 %r5, %r3, %r1, %c1;
    selp.u64 %r5, %r4, %r5, %c2;
    mov.u64 out, %r5;
    ret;
}
.func (.reg .u64 out) fpm_n (.reg .u64 base)
{
    .reg .u64 %r0, %r1, %r2;
    .reg .pred %c;
    mov.u64 %r0, base;
    ld.const.u64 %r2, [N];
    sub.u64 %r2, %r2, 2;
    mov.u64 out, 1;
lp:
    and.b64 %r1, %r2, 1;
    setp.eq.u64 %c, %r1, 0;
    @%c bra nm;
    call (out), mod_mul_n, (%r0, out);
nm:
    call (%r0), mod_mul_n, (%r0, %r0);
    setp.ne.u64 %c, %r2, 0;
    shr.u64 %r2, %r2, 1;
    @%c bra lp;
    ret;
}
.func (.reg .u64 ox, .reg .u64 oy) pta (.reg .u64 x, .reg .u64 y, .reg .u64 z)
{
    .reg .u64 %r0, %r1;
    call (%r1), mod_mul_p, (z, z);
    call (%r0), fpm_p, (%r1);
    call (ox), mod_mul_p, (%r0, x);
    call (%r0), mod_mul_p, (%r1, z);
    call (%r0), fpm_p, (%r0);
    call (oy), mod_mul_p, (%r0, y);
    ret;
}
.func (.reg .u64 ox, .reg .u64 oy, .reg .u64 oz) pd (.reg .u64 x, .reg .u64 y, .reg .u64 z)
{
    .reg .u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7;
    call (%r3), mod_mul_p, (y, 2);
    call (oz), mod_mul_p, (z, %r3);
    call (%r3), mod_mul_p, (z, z);
    call (%r3), mod_mul_p, (%r3, %r3);
    ld.const.u64 %r1, [A];
    call (%r5), mod_mul_p, (%r1, %r3);
    call (%r3), mod_mul_p, (x, x);
    call (%r3), mod_mul_p, (%r3, 3);
    call (%r6), mod_add_p, (%r3, %r5);
    call (%r7), mod_mul_p, (y, y);
    call (%r3), mod_mul_p, (x, 4);
    call (%r0), mod_mul_p, (%r3, %r7);
    call (%r3), mod_mul_p, (%r7, %r7);
    call (%r2), mod_mul_p, (%r3, 8);
    call (%r4), mod_mul_p, (%r0, 2);
    call (%r3), mod_mul_p, (%r6, %r6);
    call (ox), mod_sub_p, (%r3, %r4);
    call (%r3), mod_sub_p, (%r0, ox);
    call (%r3), mod_mul_p, (%r3, %r6);
    call (oy), mod_sub_p, (%r3, %r2);
    ret;
}
.func (.reg .u64 ox, .reg .u64 oy, .reg .u64 oz) pa (.reg .u64 x1, .reg .u64 y1, .reg .u64 z1, .reg .u64 x2, .reg .u64 y2, .reg .u64 z2)
{
    .reg .u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7;
    call (%r3), mod_mul_p, (z2, z2);
    call (%r5), mod_mul_p, (x1, %r3);
    call (%r3), mod_mul_p, (y1, %r3);
    call (%r6), mod_mul_p, (z2, %r3);
    call (%r7), mod_mul_p, (z1, z1);
    call (%r3), mod_mul_p, (x2, %r7);
    call (%r2), mod_sub_p, (%r3, %r5);
    call (%r4), mod_mul_p, (%r2, %r2);
    call (%r0), mod_mul_p, (%r2, %r4);
    call (%r3), mod_mul_p, (z1, %r2);
    call (oz), mod_mul_p, (z2, %r3);
    call (%r2), mod_mul_p, (%r5, %r4);
    call (%r3), mod_mul_p, (%r2, 2);
    call (%r1), mod_add_p, (%r3, %r0);
    call (%r3), mod_mul_p, (y2, %r7);
    call (%r3), mod_mul_p, (z1, %r3);
    call (%r5), mod_sub_p, (%r3, %r6);
    call (%r3), mod_mul_p, (%r5, %r5);
    call (ox), mod_sub_p, (%r3, %r1);
    call (%r3), mod_sub_p, (%r2, ox);
    call (%r5), mod_mul_p, (%r3, %r5);
    call (%r3), mod_mul_p, (%r6, %r0);
    call (oy), mod_sub_p, (%r5, %r3);
    ret;
}
.func (.reg .u64 ox, .reg .u64 oy) sp_mul (.reg .u64 x, .reg .u64 y, .param .u64 n)
{
    .reg .u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9, %r10;
    .reg .u32 %R0;
    .reg .pred %c;
    ld.param.u64 %r1, [n];
    bfind.u64 %R0, %r1;
    shl.b64 %r4, %r1, 1;
    xor.b64 %r4, %r4, %r1;
    call (%r2, %r5, %r6), pd, (x, y, 1);
    mov.u64 %r7, x;
    mov.u64 %r8, y;
    mov.u64 %r9, 1;
lp:
    bfe.u64 %r10, %r4, %R0, 1;
    setp.eq.u64 %c, %r10, 1;
    mov.u64 %r3, %r2;
    selp.u64 %r2, %r7, %r2, %c;
    selp.u64 %r7, %r3, %r7, %c;
    mov.u64 %r3, %r5;
    selp.u64 %r5, %r8, %r5, %c;
    selp.u64 %r8, %r3, %r8, %c;
    mov.u64 %r3, %r6;
    selp.u64 %r6, %r9, %r6, %c;
    selp.u64 %r9, %r3, %r9, %c;
    call (%r7, %r8, %r9), pa, (%r2, %r5, %r6, %r7, %r8, %r9);
    call (%r2, %r5, %r6), pd, (%r2, %r5, %r6);
    sub.u32 %R0, %R0, 1;
    setp.ne.u32 %c, %R0, 0;
    @%c bra lp;
    and.b64 %r4, %r4, 1;
    setp.eq.u64 %c, %r4, 1;
    selp.u64 %r2, %r7, %r2, %c;
    selp.u64 %r5, %r8, %r5, %c;
    selp.u64 %r6, %r9, %r6, %c;
    call (ox, oy), pta, (%r2, %r5, %r6);
    ret;
}
.entry init_M()
{
    .reg .u64 %a1, %a2, %a3, %a4;
    .reg .u64 %r1, %r2, %r3, %r4;
    .reg .u32 %R1, %R2, %R3, %R4;
    .reg .pred %c;
    ld.const.u64 %a1, [G_x];
    ld.const.u64 %a2, [G_y];
    ld.const.u64 %a3, [T];
    ld.const.u64 %a4, [M];
    mov.u32 %R2, %tid.y;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %tid.x;
    mad.lo.u32 %R1, %R2, %R3, %R4;
    mov.u32 %R2, %ctaid.x;
    mov.u32 %R3, %nctaid.y;
    mov.u32 %R4, %ctaid.y;
    mad.lo.u32 %R2, %R2, %R3, %R4;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %ntid.y;
    mul.lo.u32 %R3, %R3, %R4;
    mul.lo.u32 %R2, %R2, %R3;
    add.u32 %R1, %R1, %R2;
    cvt.u64.u32 %r1, %R1;//线程id
    mad.lo.u64 %r3, %r1, 0x18, %a4;
    ld.global.cg.u64 %r4, [%r3];
    rem.u64 %r4, %r4, %a3;
    add.u64 %r4, %r4, 1;
    setp.ne.u64 %c, %r4, 1;
    @%c bra n1;
    st.global.u64 [%r3], %a1;
    add.u64 %r3, %r3, 8;
    st.global.u64 [%r3], %a2;
    add.u64 %r3, %r3, 8;
    st.global.u64 [%r3], 1;
    ret;
n1:
    call (%r1, %r2), sp_mul, (%a1, %a2, %r4);
    st.global.u64 [%r3], %r1;
    add.u64 %r3, %r3, 8;
    st.global.u64 [%r3], %r2;
    add.u64 %r3, %r3, 8;
    st.global.u64 [%r3], %r4;
    ret;
}
.entry kangaroo(.param .u64 ret)
{
    .reg .u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9, %r10, %r11, %r12, %r13, %r14, %r15;
    .reg .pred %c0, %c1;
    .reg .u32 %R1, %R2, %R3, %R4;
    mov.u32 %R2, %tid.y;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %tid.x;
    mad.lo.u32 %R1, %R2, %R3, %R4;
    mov.u32 %R2, %ctaid.x;
    mov.u32 %R3, %nctaid.y;
    mov.u32 %R4, %ctaid.y;
    mad.lo.u32 %R2, %R2, %R3, %R4;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %ntid.y;
    mul.lo.u32 %R3, %R3, %R4;
    mul.lo.u32 %R2, %R2, %R3;
    add.u32 %R1, %R1, %R2;
    cvt.u64.u32 %r3, %R1;//线程id
    mad.lo.u64 %r0, %r3, 0x68, %r0;
    ld.const.u64 %r1, [T];
    ld.const.u64 %r2, [K];
    ld.const.u64 %r4, [M];
    ld.const.u64 %r7, [N];
    ld.const.u64 %r8, [Q];
    mov.u64 %r5, %r3;
    shr.b64 %r5, %r5, 3;
    mad.lo.u64 %r9, %r5, %r8, %r8;
    mad.lo.u64 %r8, %r5, %r8, 0;
    mov.u64 %r11, %r9;
    mad.lo.u64 %r3, %r3, %r2, 0;
    ld.const.u64 %r13, [G_x];
    ld.const.u64 %r12, [G_y];
    call (%r13, %r12), sp_mul, (%r13, %r12, %r9);
    mov.u64 %r5, 0;
lp1:
    mov.u64 %r14, %r13;
    xor.b64 %r14, %r14, %r12;
    rem.u64 %r14, %r14, %r2;
    add.u64 %r14, %r3, %r14;
    mad.lo.u64 %r14, %r14, 0x18, %r4;
    ld.global.cg.u64 %r0, [%r14];
    add.u64 %r14, %r14, 8;
    ld.global.cg.u64 %r10, [%r14];
    add.u64 %r14, %r14, 8;
    ld.global.cg.u64 %r14, [%r14];
    call (%r13, %r12, %r0), pa, (%r13, %r12, 1, %r0, %r10, 1);
    call (%r13, %r12), pta, (%r13, %r12, %r0);
    add.u64 %r11, %r11, %r14;
    add.u64 %r5, %r5, 1;
    setp.ne.u64 %c0, %r5, %r1;
    @%c0 bra lp1;
    mov.u64 %r15, %r12;
    ld.const.u64 %r10, [p_y];
    ld.const.u64 %r12, [p_x];
    mov.u64 %r5, 0;
lp2:
    xor.b64 %r6, %r12, %r10;
    rem.u64 %r6, %r6, %r2;
    add.u64 %r14, %r3, %r6;
    mad.lo.u64 %r14, %r14, 0x18, %r4;
    ld.global.cg.u64 %r0, [%r14];
    add.u64 %r14, %r14, 8;
    ld.global.cg.u64 %r1, [%r14];
    add.u64 %r14, %r14, 8;
    ld.global.cg.u64 %r14, [%r14];
    add.u64 %r5, %r5, %r14;
    call (%r12, %r10, %r0), pa, (%r12, %r10, 1, %r0, %r1, 1);
    call (%r12, %r10), pta, (%r12, %r10, %r0);
    sub.u64 %r6, %r11, %r5;
    setp.lt.u64 %c0, %r11, %r5;
    setp.lt.u64 %c1, %r6, %r8;
    or.pred %c0, %c0, %c1;
    @%c0 bra z;
    setp.ge.u64 %c0, %r6, %r9;
    setp.ne.u64 %c1, %r12, %r13;
    or.pred %c0, %c0, %c1;
    setp.ne.u64 %c1, %r10, %r15;
    or.pred %c0, %c0, %c1;
    @%c0 bra lp2;
    ld.param.u64 %r0, [ret];
    rem.u64 %r6, %r6, %r7;
    atom.or.b64 %r0, [%r0], %r6;
z:
    ret;
}
.entry init_rho()
{
    .reg .u64 %a1, %a2, %a3, %a4, %a5, %a6;
    .reg .u64 %r1, %r2, %r3, %r4, %r5, %r6;
    .reg .u32 %R1, %R2, %R3, %R4;
    ld.const.u64 %a1, [G_x];
    ld.const.u64 %a2, [G_y];
    ld.const.u64 %a3, [p_x];
    ld.const.u64 %a4, [p_y];
    ld.const.u64 %a5, [m_x];
    ld.const.u64 %a6, [m_y];
    mov.u32 %R2, %tid.y;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %tid.x;
    mad.lo.u32 %R1, %R2, %R3, %R4;
    mov.u32 %R2, %ctaid.x;
    mov.u32 %R3, %nctaid.y;
    mov.u32 %R4, %ctaid.y;
    mad.lo.u32 %R2, %R2, %R3, %R4;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %ntid.y;
    mul.lo.u32 %R3, %R3, %R4;
    mul.lo.u32 %R2, %R2, %R3;
    add.u32 %R1, %R1, %R2;
    cvt.u64.u32 %r1, %R1;//线程id
    mad.lo.u64 %r2, %r1, 8, %a5;
    ld.global.u64 %r2, [%r2];
    call (%r2, %r3), sp_mul, (%a1, %a2, %r2);
    mad.lo.u64 %r4, %r1, 8, %a6;
    ld.global.u64 %r4, [%r4];
    call (%r4, %r5), sp_mul, (%a3, %a4, %r4);
    call (%r2, %r3, %r6), pa, (%r2, %r3, 1, %r4, %r5, 1);
    call (%r2, %r3), pta, (%r2, %r3, %r6);
    mad.lo.u64 %r4, %r1, 8, %a5;
    st.global.u64 [%r4], %r2;
    mad.lo.u64 %r5, %r1, 8, %a6;
    st.global.u64 [%r5], %r3;
    ret;
}
.entry rho(.param .u64 ret)
{
    .reg .u64 %a0, %a1, %a2, %a3, %a4, %a5, %a6;
    .reg .u64 %r0, %r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8;
    .reg .pred %c0, %c1;
    .reg .u32 %R1, %R2, %R3, %R4;
    .reg .u64 %keys, %vas, %vbs, %m_x, %m_y, %arr1, %arr2;
    ld.const.u64 %a0, [T];
    ld.const.u64 %a1, [N];
    ld.const.u64 %a2, [G_x];
    ld.const.u64 %a3, [G_y];
    ld.const.u64 %a4, [p_x];
    ld.const.u64 %a5, [p_y];
    ld.const.u64 %a6, [MSK];
    mov.u32 %R2, %tid.y;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %tid.x;
    mad.lo.u32 %R1, %R2, %R3, %R4;
    mov.u32 %R2, %ctaid.x;
    mov.u32 %R3, %nctaid.y;
    mov.u32 %R4, %ctaid.y;
    mad.lo.u32 %R2, %R2, %R3, %R4;
    mov.u32 %R3, %ntid.x;
    mov.u32 %R4, %ntid.y;
    mul.lo.u32 %R3, %R3, %R4;
    mul.lo.u32 %R2, %R2, %R3;
    add.u32 %R1, %R1, %R2;
    cvt.u64.u32 %r1, %R1;//线程id
    ld.const.u64 %keys, [keys];
    ld.const.u64 %vas, [vas];
    ld.const.u64 %vbs, [vbs];
    ld.const.u64 %m_x, [m_x];
    ld.const.u64 %m_y, [m_y];
    mov.u64 %arr1, arr1;
    mov.u64 %arr2, arr2;
    mad.lo.u64 %keys, %r1, 0x20000, %keys;//M_S*8==0x20000
    mad.lo.u64 %vas, %r1, 0x20000, %vas;
    mad.lo.u64 %vbs, %r1, 0x20000, %vbs;
    and.b64 %r2, %r1, 0xf00;
    add.u64 %m_x, %m_x, %r2;
    add.u64 %m_y, %m_y, %r2;
    add.u64 %arr1, %arr1, %r2;
    add.u64 %arr2, %arr2, %r2;
    shr.b64 %r2, %r1, 4;
    ld.const.u64 %r2, arr1[%r2];
    and.b64 %r3, %r1, 0x3f;
    ld.const.u64 %r3, arr2[%r3];
    add.u64 %r1, %r1, %r3;
    rem.u64 %r1, %r1, %a1;
    add.u64 %r1, %r1, 2;
    call (%r3, %r4), sp_mul, (%a2, %a3, %r1);
    rem.u64 %r1, %r1, %a1;
    mov.u64 %r2, 0;
    mov.u64 %r0, 0;
lp:
    xor.b64 %r8, %r3, %r4;
    rem.u64 %r5, %r8, 0x20;
    mad.lo.u64 %r6, %r5, 8, %m_x;
    ld.global.u64 %r6, [%r6];
    mad.lo.u64 %r7, %r5, 8, %m_y;
    ld.global.u64 %r7, [%r7];
    call (%r3, %r4, %r6), pa, (%r3, %r4, 1, %r6, %r7, 1);
    call (%r3, %r4), pta, (%r3, %r4, %r6);
    mad.lo.u64 %r6, %r5, 8, %arr1;
    ld.const.u64 %r6, [%r6]; 
    call (%r1), mod_add_n, (%r1, %r6);
    mad.lo.u64 %r6, %r5, 8, %arr2;
    ld.const.u64 %r6, [%r6];
    call (%r2), mod_add_n, (%r2, %r6);
    and.b64 %r5, %r8, 0x3fff;//M_S-1==0x3fff
    mad.lo.u64 %r6, %r5, 8, %keys;
    ld.global.u64 %r6, [%r6];
    setp.eq.u64 %c0, %r6, %r8;
    @%c0 bra fnd;
nf:
    and.b64 %r7, %r0, %a6;
    setp.ne.u64 %c0, %r7, 0;
    @%c0 bra ne;
    mad.lo.u64 %r6, %r5, 8, %keys;
    st.global.u64 [%r6], %r8;
    mad.lo.u64 %r6, %r5, 8, %vas;
    st.global.u64 [%r6], %r1;
    mad.lo.u64 %r6, %r5, 8, %vbs;
    st.global.u64 [%r6], %r2;
ne:
    add.u64 %r0, %r0, 1;
    setp.ne.u64 %c0, %r0, %a0;
    @%c0 bra lp;
    ret;
fnd:
    mad.lo.u64 %r6, %r5, 8, %vbs;
    ld.global.u64 %r6, [%r6];
    setp.eq.u64 %c0, %r2, %r6;
    @%c0 bra ed;
    mad.lo.u64 %r5, %r5, 8, %vas;
    ld.global.u64 %r5, [%r5];
    call (%r5), mod_sub_n, (%r5, %r1);
    call (%r2), mod_sub_n, (%r2, %r6);
    call (%r2), fpm_n, (%r2);
    call (%r5), mod_mul_n, (%r5, %r2);
    setp.le.u64 %c0, %r5, 1;
    @%c0 bra ed;
    call (%r1, %r2), sp_mul, (%a2, %a3, %r5);
    setp.ne.u64 %c0, %a4, %r1;
    setp.ne.u64 %c1, %a5, %r2;
    or.pred %c0, %c0, %c1;
    selp.u64 %r5, 0, %r5, %c0;
    ld.param.u64 %r0, [ret];
    atom.or.b64 %r0, [%r0], %r5;
ed:
    ret;
}
```
观察发现,并行rho性能表现较好,但是仍未与cpu端单线程rho拉开差距,还需进一步优化.

并行kangaroo算法的效率约为并行rho的1/8,这两种并行算法效率抖动大小相近.

利用此并行rho算法,题目的所有所需攻击(数次2^64级)大约耗时4小时.

参考文献:官方文档 <mark class="encrypted">770K9s2c8@1M7s2y4Q4x3@1q4Q4x3V1k6Q4x3V1k6V1L8$3y4K6i4K6u0W2L8Y4k6A6k6r3W2S2i4K6u0W2j5$3!0E0i4K6u0r3j5%4g2V1j5g2)9J5c8Y4m8S2M7X3q4D9L8r3g2D9i4K6u0V1N6r3S2J5k6h3q4V1i4K6u0V1k6i4S2W2j5%4g2@1K9h3!0F1i4K6u0r3K9h3&6V1k6i4S2Q4x3X3g2Z5N6r3#2D9</mark>

登录后可查看完整内容

冰与火的战歌：Windows内核攻防实战高级班！从零到实战，融合AI与Windows内核攻防全技术栈，打造具备自动化能力的内核开发高手。

最后于 2026-5-6 17:15 被n00bzx编辑，原因：

#CrackMe #Reverse

收藏・1

点赞・1

打赏

最新回复 (1)
n00bzx 雪币： 2703 活跃值： (3839) 能力值： ( LV12，RANK：286 ) 在线值：发帖 7 回帖 52 粉丝 10 关注私信	n00bzx 2 2 楼好耶!看雪有滑稽哦! 2026-5-7 11:55 0
	游客登录 \| 注册方可回帖回帖表情雪币赚取及消费高级回复

n00bzx

发帖

回帖

286

RANK

关注

私信

他的文章

关于我们

联系我们

企业服务

看雪公众号

最新回复 (1)
n00bzx 雪币： 2703 活跃值： (3839) 能力值： ( LV12，RANK：286 ) 在线值：发帖 7 回帖 52 粉丝 10 关注私信	n00bzx 2 2 楼好耶!看雪有滑稽哦! 2026-5-7 11:55 0
	游客登录 \| 注册方可回帖回帖表情雪币赚取及消费高级回复