[<prev] [next>] [<thread-prev] [thread-next>] [day] [month] [year] [list]
Date: Fri, 16 Aug 2013 02:16:44 +0200
From: Dániel Bali <balijanosdaniel@...il.com>
To: john-dev@...ts.openwall.com
Subject: Re: Daniel's weekly report #9
Hello!
Yesterday and today I was looking at the ISA and tried to get patching to
work with the binary generated from bcrypt-opencl.
I rewrote the code that initializes the buffers for the to-be-patched
bcrypt kernel, but it doesn't work.
Here is the code:
https://gist.github.com/balidani/22409a51df592c4b3b93
It is mostly the same as
binary_gen.c<https://github.com/balidani/gcnasm/blob/master/tools/opencl_kernel_binary_gen/binary_gen.c>from
the repository, except it sets more kernel arguments. The definition
of the blowfish kernel function is:
__kernel void blowfish( constant uint *salt
__attribute__((max_constant_size(16))),
constant uint *P_box __attribute__((max_constant_size(72))),
__global uint *BF_out,
__global uint *BF_current_S,
__global uint *BF_current_P_global,
uint rounds,
constant uint *S_box __attribute__((max_constant_size(4096))) )
{ ... }
This means we have to pass 2 constant buffers, then 3 normal ones, an
integer and then a final constant buffer. This is what the code I linked to
does. 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;
}
It did not work. The result was an empty buffer.
I don't have any ideas about proceeding here.
Does anybody have one?
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.
