Follow @Openwall on Twitter for new release announcements and other news
[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Wed, 21 Aug 2013 00:35:33 +0200
From: Dániel Bali <>
Subject: Re: Daniel's weekly report #10


2013/8/20 Lukas Odzioba <>

> 2013/8/20 Dániel Bali <>:
> > 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 <>:
> > 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:
(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.


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.