Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<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

Your e-mail address:

Powered by Openwall GNU/*/Linux - Powered by OpenVZ