[<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
Confused about mailing lists and their use?
Read about mailing lists on Wikipedia
and check out these
guidelines on proper formatting of your messages.
