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


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:

It is mostly the same as
the repository, except it sets more kernel arguments. The definition
of the blowfish kernel function is:

__kernel void blowfish( constant uint *salt
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?


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.