# include  <stdio.h> 

int main ( );
__global__ void collatz_steps ( int *steps_gpu );

/******************************************************************************/

int main ( ) 

/******************************************************************************/
/*
  Purpose:

    COLLATZ finds the longest Collatz sequence between 1 and N.

  Modified:

    07 March 2017

  Author:

    John Burkardt

  Parameters:

    Local, int *STEPS_CPU, will contain the length of each Collatz
    sequence from 1 to STEPS_NUM.

    Local, int STEPS_NUM, the highest integer to check.
*/
{
  int i;
  int i_max;
  int *steps_cpu;
  int *steps_gpu; 
  int steps_num;
  int steps_size;

  printf ( "\n" );
  printf ( "COLLATZ:\n" );
  printf ( "  CUDA/C version\n" );
//
// Allocate space on the GPU.
//
  steps_num = 100;
  steps_size = steps_num * sizeof ( int );
  cudaMalloc ( ( void** ) &steps_gpu, steps_size ); 
//
// Launch the kernel on the GPU, using 25 block and 4 threads per block.
//
  int blocks = 25;
  int threads = 4;
  collatz_steps <<< blocks, threads >>> ( steps_gpu ); 
//
// Copy steps_cpu <= steps_gpu from device to host.
//
  steps_cpu = ( int * ) malloc ( steps_size );
  cudaMemcpy ( steps_cpu, steps_gpu, steps_size, cudaMemcpyDeviceToHost );
//
//  Find the maximum sequence length.
//
  i_max = 0;
  for ( i = 0; i < steps_num; i++ ) 
  {
    if ( steps_cpu[i_max] < steps_cpu[i] )
    {
      i_max = i;
    }
  }
  printf ( "\n" );
  printf ( "  In the range [1,%d], the longest Collatz sequence\n", steps_num );
  printf ( "  occurs for N = %d, of length %d.\n", i_max + 1, steps_cpu[i_max] );
// 
// Free memory on the GPU.
//
  cudaFree ( steps_gpu ); 
//
//  Free memory on the host.
//
  free ( steps_cpu ); 
//
//  Terminate.
//
  printf ( "\n" );
  printf ( "COLLATZ:\n" );
  printf ( "  Normal end of execution.\n" );

  return 0;
}
/******************************************************************************/

__global__ void collatz_steps ( int *steps_gpu )

/******************************************************************************/
/*
  Purpose:

    COLLATZ_STEPS computes the length of a particular Collatz sequence.

  Discussion:

    We assume the kernel is called with scalar values of BLOCKS and THREADS.
    Task K checks integer N+1.

  Parameters:

    Output, int *STEPS_GPU, the array of Collatz sequence lengths.
    This function will set the K-th value of this array.
*/
{
  int k;
  int n;
  int s;

  k = threadIdx.x + blockDim.x * blockIdx.x; 
  n = k + 1;
  s = 0;
  
  while ( 1 < n )
  {
    if ( ( n % 2 ) == 0 )
    {
      n = n / 2;
    }
    else
    {
      n = 3 * n + 1;
    }
    s = s + 1;
  }

  steps_gpu[k] = s;

  return;
}