Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [day] [month] [year] [list]
Date: Thu, 27 Jun 2013 17:13:08 +0200
From: Dániel Bali <balijanosdaniel@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: Daniel's weekly report #2

Hello!


2013/6/27 Solar Designer <solar@...nwall.com>

> On Mon, Jun 24, 2013 at 03:26:47PM +0200, D?niel Bali wrote:
>
> > 2. Removed the code duplication from the keccak256 format with the help
> of
> > Lukas
>
> Is it working reliably, at least on both GPUs in bull?
>

I ran it a couple of times both on bull and on my laptop and it worked.


> > Priorities:
> > 1. Get the assembler working on a small subset of op codes
> > 2. Try patching the "dummy kernel" that contains NOPs with working
> > microcode that changes the behavior of the kernel.
> > 3. Try to figure out how the kernel's initialization works
> > 4. Expand the subset of supported instructions
>
> How much progress at these do you expect to make this week?
>

Right now the assembler supports SOP2 type instructions with SGPRs and
literal values (plus inline constants) as the operands. I should add
support for other operand types soon (probably tonight).

Lukas asked me to look at kernel initialization so I will definitely spend
time on that as well.

How are you going to produce that dummy kernel specifically with NOPs?
> Request some insane code alignment via an OpenCL directive maybe?  Or
> generate a kernel with other than NOPs (with some dummy non-NOP code)
> and overwrite its code with NOPs using your would-be assembler?
>

When I was applying for the GSoC project I made a dummy kernel with NOPs
like that (the second version you list).
I created non-NOP code and then changed all the instructions to NOPs by
"hand" (some python script).
Creating long non-NOP code wasn't trivial because OpenCL has pretty good
optimization.

With Lukas we decided to work with only a simple kernel signature (is it
called that?) for now, which is

 __kernel void doSomething(__global uint *input, __constant uint *salt,
__gloal uint *output)

Regards,
Daniel

Content of type "text/html" skipped

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.