# include <stdio.h> 
# include <stdlib.h>
# include <time.h>

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

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

int main ( ) 

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

    COLLATZ finds the longest Collatz sequence between 1 and N.

  Modified:

    11 May 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.
*/
{
  clock_t clock1;
  clock_t clock2;
  int i;
  int i_max;
  int *steps_cpu;
  int *steps_gpu; 
  int steps_num = 10000000;
  int steps_size;
  double value;

  printf ( "\n" );
  printf ( "COLLATZ:\n" );
  printf ( "  CUDA version\n" );
//
// Allocate space on the GPU.
//
  steps_size = steps_num * sizeof ( int );
  cudaMalloc ( ( void** ) &steps_gpu, steps_size ); 
//
//  Allocate space on the CPU.
//
  steps_cpu = ( int * ) malloc ( steps_size );
//
// Launch the kernel on the GPU, using 1 block and STEPS_NUM threads.
//
  clock1 = clock ( );
  collatz_steps <<< 1, steps_num >>> ( steps_gpu ); 
//
// Copy steps_cpu <= steps_gpu from device to host.
//
  cudaMemcpy ( steps_cpu, steps_gpu, steps_size, cudaMemcpyDeviceToHost );
//
//  Wait until the result is returned before measuring time.
//
  clock2 = clock ( );

  value = ( double ) ( clock2 - clock1 ) / ( double ) CLOCKS_PER_SEC;
  printf ( "\n" );
  printf ( "  Computation required %g seconds\n", value );
//
//  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:

    The value of I, the thread ID given to this function by the GPU,
    is used to specify N=I+1, the integer to be checked.

  Parameters:

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

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

  steps_gpu[i] = s;

  return;
}