Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Wed, 21 Aug 2013 00:35:33 +0200
From: Dániel Bali <balijanosdaniel@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: Daniel's weekly report #10

Hello!


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

> 2013/8/20 Dániel Bali <balijanosdaniel@...il.com>:
> > 1b. Changed the OpenCL code to a minimalistic example. Changing one bit
> of
> > this code results in a lack of LDS usage.
>
> I don't buy it, what did you try to change?
>

Here it is. There are 2 versions of the last line.

__kernel void sample(
    __global int *in,
    __global int *in2,
    __global int *out)
{
    __local int w[128];
    int gid = get_global_id(0);

    out[gid] = in[gid] + w[0]; // generated binary uses the LDS
    out[gid] = in[gid]; // generated binary does not use the LDS
}


>
> > 3. Tried to write a dummy that uses 64-bit unsigned integer buffers, but
> > something went wrong with the patching. I will work on that.
>
> What's the current status of that?
>

2013/8/20 Dániel Bali <balijanosdaniel@...il.com>:
> > 1. Get the 64-bit binary patching working
> > 1a. Get ds_add_u64 working (there is not much documentation on it)
> What approach do you use to accomplish that?
> Have you tried to extend my opencl code to use 64bit integers?


Yes, this is what I tried, but it does not work.
I changed the binary_gen.c source to pass 64-bit integers to the OpenCL
kernel and then changed your code to use ulong. Something goes wrong when
patching the 64-bit dummy. Simply loading a kernel with the "binary_gen"
binary works perfectly, but when it comes to patching it fails.

I'll give an update on this when I know where the error is.
This is what the 64-bit binary_gen source looks like:
https://gist.github.com/balidani/df5a4eff8618ea94efd2
(Only slight differences)

Trying to load a short OpenCL kernel with this works, but generating the
ISA and then patching it into a dummy does not.


> > 2. Continue trying to figure out the LDS problem
> I am not sure about what LDS problem you're talking now?
> Is it related to non documented sections of generated binaries?


Yes. Basically I have no idea where LDS is "turned on" in the ELF.
Something must turn it on, because there are differences in the LDS and
non-LDS binaries, but looking at the differences in the ATI CAL notes
doesn't show anything promising.

I could upload these binaries somewhere if you are interested.

Regards,
Daniel

Content of type "text/html" skipped

Powered by blists - more mailing lists

Your e-mail address:

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