Edit Question - KnowledgeHub

Title Body

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. 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. I appreciate if somebody can give me an idea how to realize efficient modular computations with CUDA ? To see how this is implemented on CUDA, I use the following code snippet: __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; } This code is not supposed to work, I just wanted to see how modular reduction is implemented on CUDA. When I disassemble this with cuobjdump --dump-sass (thanks njuffa for advice!), I see the following: /*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; 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

Tags (comma-separated) Cancel