8 Replies Latest reply on Aug 4, 2014 2:39 AM by pinform

    rotate() broken

    atom

      Here is an OpenCL Kernel which fails to compute some of the values correctly. I've tested this on Ubuntu 12.04 with Catalyst 13.4 but it fails on older versions, too. Reproduces on all my tested cards. Problem is located somewhere in the builtin rotate().

       

      /*  This kernel returns the following values:

        

      0: 0x01234567 = 0x01234567

      1: 0x02468ace = 0x02468ace

      2: 0x048d159c = 0x048d159c

      3: 0x091a2b38 = 0x091a2b38

      4: 0x12345670 = 0x12345670

      5: 0x2468ace0 = 0x2468ace0

      6: 0x48d159c0 = 0x48d159c0

      7: 0x91a2b380 = 0x91a2b380

      8: 0x67012345 = 0x23456701 <- ERROR

      9: 0x468ace02 = 0x468ace02

      10: 0x8d159c04 = 0x8d159c04

      11: 0x1a2b3809 = 0x1a2b3809

      12: 0x34567012 = 0x34567012

      13: 0x68ace024 = 0x68ace024

      14: 0xd159c048 = 0xd159c048

      15: 0xa2b38091 = 0xa2b38091

      16: 0x45670123 = 0x45670123

      17: 0x8ace0246 = 0x8ace0246

      18: 0x159c048d = 0x159c048d

      19: 0x2b38091a = 0x2b38091a

      20: 0x56701234 = 0x56701234

      21: 0xace02468 = 0xace02468

      22: 0x59c048d1 = 0x59c048d1

      23: 0xb38091a2 = 0xb38091a2

      24: 0x23456701 = 0x67012345 <- ERROR

      25: 0xce02468a = 0xce02468a

      26: 0x9c048d15 = 0x9c048d15

      27: 0x38091a2b = 0x38091a2b

      28: 0x70123456 = 0x70123456

      29: 0xe02468ac = 0xe02468ac

      30: 0xc048d159 = 0xc048d159

      31: 0x8091a2b3 = 0x8091a2b3

       

      uint ROTATE_LEFT (const uint a, const uint n)

      {

        return (a << n) | (a >> (32 - n));

      }

       

      __kernel void oclHashcat_RipeMD160 (__global uint *results)

      {

        const int gid = get_global_id (0);

       

        if (gid != 0) return; // this is just to avoid race conditions while writing results

       

        const uint v = 0x01234567;

       

        for (uint i = 0; i < 32; i++) results[ 0 + i] = rotate      (v, i);

        for (uint i = 0; i < 32; i++) results[32 + i] = ROTATE_LEFT (v, i);

      }

       

      Regards,

      Jens

        • Re: rotate() broken
          kylesiefring

          It's not entirely clear from atom's post, but 0x23456701 is the desired output for results[8]. Looking closely, I can see that the results for 8 and 24 are switched. If we were rotating right, then the results for 8 and 24 would be correct. 16 is correct if we are rotating right or left. Of course, we are rotating left and not right. I'm guessing some compiler optimization was made for increments of 8 and the coder mistakenly put the wrong direction in.

           

          I try doing something to make the loop not unroll (it might not be anyway), but Visual Studio is crashing on me whenever I press enter.

          • Re: rotate() broken
            himanshu.gautam

            Thanks for reporting it. I will try to see if it happens at my end too.