Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Fri, 16 Aug 2013 11:24:50 +0200
From: Dániel Bali <balijanosdaniel@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: Re: Daniel's weekly report #9

Hello!

2013/8/16 Lukas Odzioba <lukas.odzioba@...il.com>

> 2013/8/16 Dániel Bali <balijanosdaniel@...il.com>:
> > I tried to patch the binary with the microcode (generated by
> > KernelAnalyzer) of the following opencl code:
> >
> > __kernel void blowfish(constant uint *buf1, constant uint *buf2, __global
> > uint *buf3, __global uint *buf4, __global uint *buf5, uint rounds,
> constant
> > uint *buf6) {
> >     buf3[0] = 111;
> >     buf4[0] = 222;
> >     buf5[0] = 333;
> > }
>
> Can you post this isa code?
>

I attached it to the mail.


>
> I am affraid you're still missing the point of this test. If using bf
> kernel to test ds instructions is hard, maybe it will be faster/easier
> to write simple opencl code that will utilize ds instructions itself,
> and will be easier to modify.
>
>
I'm unable to produce any ISA with the KernelAnalyzer that has ds_
instructions in them.
Even when using the local id, like this:

__kernel void sample(
    __global uint *input,
    __local uint *buf,
    __global uint *output) {

        int gid = get_global_id(0);
        int lid = get_local_id(0);

        output[gid + lid] = input[gid] + input[gid + lid];
}

Regards,
Daniel

Content of type "text/html" skipped

Download attachment "bf.isa" of type "application/octet-stream" (4064 bytes)

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.