AnsweredAssumed Answered

Problem with bit operations with openCL

Question asked by snowsp on Dec 30, 2013
Latest reply on Mar 13, 2014 by amd_support

Hi, I did some tests on an HD7970 card recently, simply running the md5 hash function in openCL on it. But I found that the result is wrong (the same code runs fine on an NVIDIA card).

It became correct if I used the cl-disable-opt option. Does anyone have similar experience? Is it a bug in the compiler? I have been using openCL 1.2.

A simplified version of the kernel code is like this:

__kernel void crack(__global unsigned char * correctPass, __global unsigned char *cudaCharSet, int wordLength, int charSetLen, unsigned int v1, unsigned int v2, unsigned int v3, unsigned int v4, __constant unsigned char *cudaBrute, __global unsigned int* result)

{

  const unsigned int a0 = 0x67452301;

  const unsigned int b0 = 0xEFCDAB89;

  const unsigned int c0 = 0x98BADCFE;

  const unsigned int d0 = 0x10325476;

  int length = 3;

  unsigned int a = 0;

  unsigned int b = 0;

  unsigned int c = 0;

  unsigned int d = 0;

 

     vals[0] |= ((unsigned int)v1) << ((0 % 4) * 8);         //(97+((int)(get_group_id(0))%26)) << ((0 % 4) * 8);

 

      vals[0] |= (97+1%26) << ((1 % 4) * 8);

 

      vals[0] |= (97+2) << ((2 % 4) * 8);

 

     vals[0] |= 0x80 << ((3 % 4) * 8);

 

 

  //Initialize hash value for this chunk:

  a = a0;

  b = b0;

  c = c0;

  d = d0;

 

 

#define S11 7

#define S12 12

#define S13 17

#define S14 22

  FF ( a, b, c, d, vals[0],  S11, 3614090360); /* 1 */

  FF ( d, a, b, c, 0,  S12, 3905402710); /* 2 */

  FF ( c, d, a, b, 0,  S13,  606105819); /* 3 */

  unsigned int inter = ~c;

  //FF ( b, c, d, a, 0,  S14, 3250441966); /* 4 */

 

  result[2]=a&~c;

 

 

  unsigned int test2 = ~c&a&0x80000000;

  result[3] = (d&c) | result[2];

}

//////////////////////////////

The value of result[3] is wrong. If I used a constant value instead of "v1" in the first assignment of vals[0], the result is correct.

Outcomes