title: modular arithmetic on the gpu
body: <p>I am working on the GPU algorithm which is supposed to do a lot of modular computations. Particularly, various operations on matrices in a finite field which in the long run
  reduce to primitive operations like: (a*b - c*d) mod m or (a*b + c) mod m where a,b,c and d are residues modulo m and m is a 32-bit prime. </p>
  
  <p>Through experimentation I learned that the performance of the algorithm is mostly limited by slow modular arithmetic because integer modulo (%) and division operations are not supported on the GPU in hardware.</p>
  
  <p>I appreciate if somebody can give me an idea how to realize efficient modular computations with CUDA ?</p>
  
  <p>To see how this is implemented on CUDA, I use the following code snippet:</p>
  
  <pre><code>__global__ void mod_kernel(unsigned *gout, const unsigned *gin) {
  
  unsigned tid = threadIdx.x;
  unsigned a = gin[tid], b = gin[tid * 2], m = gin[tid * 3];
  
  typedef unsigned long long u64;
  
  __syncthreads();
  unsigned r = (unsigned)(((u64)a * (u64)b) % m);
  __syncthreads();
  gout[tid] = r;
  }
  </code></pre>
  
  <p>This code is not supposed to work, I just wanted to see how modular reduction is
  implemented on CUDA. </p>
  
  <p>When I disassemble this with cuobjdump --dump-sass (thanks njuffa for advice!), I see the following:</p>
  
  <pre><code>/*0098*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
  /*00a0*/     /*0x1c315c4350000000*/     IMUL.U32.U32.HI R5, R3, R7;
  /*00a8*/     /*0x1c311c0350000000*/     IMUL.U32.U32 R4, R3, R7;
  /*00b0*/     /*0xfc01dde428000000*/     MOV R7, RZ;
  /*00b8*/     /*0xe001000750000000*/     CAL 0xf8;
  /*00c0*/     /*0x00000007d0000000*/     BPT.DRAIN 0x0;
  /*00c8*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
  </code></pre>
  
  <p>Note that between the two calls to bar.red.popc there is a call to 0xf8 procedure which implements some sophisticated algorithm (about 50 instructions or even more). Not surpising that mod (%) operation is slow</p>