Permuting a vector

I’m working on a problem where I need to make a lot of permutations on a lot of vectors, for random permutation t-test, how should this be implemented in CUDA?

My vectors are for example of the length 100 samples and I want to do 10 000 permutations (and t-tests) on each vector.

I am NOT an expert in CUDA, just a newbie. But here is how I do it. Also, I do not check in here often, so if you have a question I may miss it.

// NEVER call this routine with iseed=0.

// Also, it will never return 0 in iseed.

// Unifrand is returned from 0 (inclusive) to 1.0 (exclusive) (for DOUBLE math. It may round up to 1.0 for FLOAT math.)

// I think the final division should be with a denominator one less than that used,

// but this is inconsequential.

device void unif ( int *iseed , float *unifrand )

{

int k ;

k = *iseed / 127773 ; // The constants in this routine are critical! Do not change them.

*iseed = 16807 * (*iseed - k * 127773) - 2836 * k ;

if (*iseed < 0)

  *iseed += 2147483647 ;

*unifrand = (*iseed - 1.0f) / 2147483647.0f ; // I think the real max is one less than this denom

}

// CUDA_shuffle_indices is a global array nthreads * n_market long

// Thread 0 is not shuffled.

global void cuda_shuffle_kernel ( int nthreads , int n_market , int *CUDA_shuffle_indices )

{

int i, j, ip, irand, itemp, *indices ;

float frand ;

ip = blockIdx.x * blockDim.x + threadIdx.x ;

if (ip >= nthreads)

  return ;

indices = CUDA_shuffle_indices + ip * n_market ;

for (i=0 ; i<n_market ; i++)

  indices[i] = i ;

if (ip) {

  irand = ip ;

  unif ( &irand , &frand ) ; // Warm up generator

  unif ( &irand , &frand ) ;

i = n_market ; // Number remaining to be shuffled

  while (i > 1) {        // While at least 2 left to shuffle

     unif ( &irand , &frand ) ;

     j = (int) (frand * i) ;

     if (j >= i)     // This is necessary only for FLOAT math.  DOUBLE guarantees j < i.

        j = i - 1 ;

     --i ;

     itemp = indices[i] ;

     indices[i] = indices[j] ;

     indices[j] = itemp ;

     }

  }

}

I’m sorry that the formatting here destroys the indentation.

Tim