Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Sun, 23 Aug 2015 11:40:35 +0300
From: Solar Designer <solar@...nwall.com>
To: john-dev@...ts.openwall.com
Subject: Re: PHC: Argon2 on GPU

On Sun, Aug 23, 2015 at 11:02:24AM +0300, Solar Designer wrote:
> Unfortunately, when we're dealing with 64-bit types, the generated PTX
> code includes extra mov's:
> 
>         {
>         .reg .b32 %dummy;
>         mov.b64         {%r15,%dummy}, %rd82;
>         }
>         {
>         .reg .b32 %dummy;
>         mov.b64         {%dummy,%r16}, %rd82;
>         }
>         shf.r.wrap.b32  %r17, %r16, %r15, 24;
>         shf.r.wrap.b32  %r18, %r15, %r16, 24;
> 
> These are simply to extract the 32-bit halves as needed for the shf
> instructions.  The mov's should be gone and proper registers
> substituted right into the shf instructions in the final ISA code -
> however, I am not sure this is what is actually happening (depends on
> how good the translator from PTX to native ISA is).

I passed the PTX code through "ptxas --gpu-name sm_35" and nvdisasm, and
it looks OK in this respect:

        /*28d0*/                   LOP.XOR R26, R248, R12;
        /*28d8*/                   LOP.XOR R36, R32, R13;
        /*28e0*/                   LOP.XOR R37, R25, R31;
        /*28e8*/                   IADD.X R239, R24, R27;
        /*28f0*/                   LDL.64 R24, [R141+0x58];
        /*28f8*/                   IADD R35.CC, R232, c[0x3][0x0];
        /*2908*/                   SHF.R.W R242, R26, 0x18, R36;
        /*2910*/                   LOP.XOR R7, R7, c[0x3][0x34];
        /*2918*/                   SHF.R.W R39, R36, 0x18, R26;

Here we can see R26 and R36 come directly from LOP.XOR, without MOV.

Also interesting is LDL.64.  I guess it loads two adjacent registers
(R24 and R25 in this example), which under the SIMT model are 32-bit
elements in two different hardware SIMD registers.

Alexander

Powered by blists - more mailing lists

Confused about mailing lists and their use? Read about mailing lists on Wikipedia and check out these guidelines on proper formatting of your messages.