Large Integers on CUDA

Accelerated Computing CUDA CUDA Programming and Performance
#1

Hi everyone,

I was wondering if there already is a large integer (big int) library for CUDA, to be run on the devices. I am writing some applications that need to deal with bigints, so any code will help me save some time.

Thanks in advance,
Renato

#2

Based on my limited understanding of big int arithmetic, it doesn’t seem likely an efficient big int library exists now. An add with carry instruction is needed to do multiprecision adds efficiently. The hardware is capable of it because I see the addc instruction mentioned in the PTX instruction set, but I don’t see it available in the header files.

Maybe NVCC is smart enough to recognize big int arithmetic & use addc?

#3

Hi,

just build your own intrinsics. It is really easy, I did it for my CUDA application, too.

__device__ unsigned int __add_cc(unsigned int a, unsigned int b)

{

  unsigned int r;

  asm("add.cc.u32 %0, %1, %2;" : "=r" (r) : "r" (a) , "r" (b));

  return r;

}

__device__ unsigned int __addc_cc(unsigned int a, unsigned int b)

{

  unsigned int r;

  asm("addc.cc.u32 %0, %1, %2;" : "=r" (r) : "r" (a) , "r" (b));

  return r;

}

__device__ unsigned int __addc(unsigned int a, unsigned int b)

{

  unsigned int r;

  asm("addc.u32 %0, %1, %2;" : "=r" (r) : "r" (a) , "r" (b));

  return r;

}

This might be incompatibly for feature CUDA versions but it works great on current versions. :)

You can access (u)mul24.hi this way, too.

Oliver

P.S. what kind of big integer functions do you need?