/*
 * Copyright 1993-2007 NVIDIA Corporation.  All rights reserved.
 *
 * NOTICE TO USER:
 *
 * This source code is subject to NVIDIA ownership rights under U.S. and
 * international Copyright laws.  Users and possessors of this source code
 * are hereby granted a nonexclusive, royalty-free license to use this code
 * in individual and commercial software.
 *
 * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
 * CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
 * IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH
 * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
 * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
 * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
 * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
 * OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
 * OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE
 * OR PERFORMANCE OF THIS SOURCE CODE.
 *
 * U.S. Government End Users.   This source code is a "commercial item" as
 * that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
 * "commercial computer  software"  and "commercial computer software
 * documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995)
 * and is provided to the U.S. Government only as a commercial end item.
 * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
 * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
 * source code with only those rights set forth herein.
 *
 * Any use of this source code in individual and commercial software must
 * include, in the user documentation and internal comments to the code,
 * the above Disclaimer and U.S. Government End Users Notice.
 */

// Compile :
// nvcc sample.cu -I(CUDA_SDK_DIR)/common/inc/
// -I(CUDA_toolkit_DIR)/include -L(CUDA_toolkit_DIR)/lib -L(CUDA_SDK_DIR)/lib
// -lcutil  -lglut
//
// Sample Program
//

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <cutil.h>

// kernel for device
__global__ void kernelDevice( float* d_a, float* d_b, float* d_c, int nx, int loadLevel){
  for(int itr=0; itr< loadLevel; itr++)
    for(int i=0;i<nx; i++){
      d_c[i] = d_a[i] * d_b[i] + 1.0;
    }
  __syncthreads();
}

////////////////////////////////////////////////////////////////////////////////
// Reference on Host
void kernelHost( float *a, float* b, float *c, int nx, int loadLevel){
  for(int itr=0; itr< loadLevel; itr++)
    for(int i=0;i<nx;i++){
      c[i] = a[i] * b[i] + 1.0;
    }
}


void runTest( int argc, char** argv);

// main function
int main( int argc, char** argv) 
{
    runTest( argc, argv);
    CUT_EXIT(argc, argv);
}

////////////////////////////////////////////////////////////////////////////////
//! Run test code for CUDA
////////////////////////////////////////////////////////////////////////////////
void runTest( int argc, char** argv) 
{
  int array_size = 100*100*100;
  int max_itr=100;
  int loadLevel=100;

  CUT_DEVICE_INIT();
  
  int mem_size = sizeof(float) * array_size;
  
  // allocate host memory
  float* h_a = (float*) malloc(mem_size);
  float* h_b = (float*) malloc(mem_size);
  
  // initalize the array with random number
  for(int i = 0; i < array_size; i++){
    h_a[i] = rand()/(float)RAND_MAX;
    h_b[i] = rand()/(float)RAND_MAX;
  }
  
  // allocate device memory
  float* d_a; float* d_b;
  CUDA_SAFE_CALL( cudaMalloc( (void**) &d_a, mem_size));
  CUDA_SAFE_CALL( cudaMalloc( (void**) &d_b, mem_size));
  
  // copy host memory to device
  CUDA_SAFE_CALL( cudaMemcpy( d_a, h_a, mem_size,
			      cudaMemcpyHostToDevice) );
  CUDA_SAFE_CALL( cudaMemcpy( d_b, h_b, mem_size,
			      cudaMemcpyHostToDevice) );
  
  // allocate device memory for result
  float* d_c;
  CUDA_SAFE_CALL( cudaMalloc( (void**) &d_c, mem_size));
  
  // setup execution parameters
  dim3  grid( 1, 1, 1);
  dim3  threads( 1, 1, 1);
  
  unsigned int timer = 0;
  CUT_SAFE_CALL( cutCreateTimer(&timer));
  CUT_SAFE_CALL( cutStartTimer(timer));
  
  // execute the kernel
  for(int i=0;i<max_itr;i++){
    if( !(i%10) ) putchar('|'); fflush(stdout); 
    putchar('.'); fflush(stdout);
    kernelDevice<<< grid, threads, 0 >>>( d_a, d_b, d_c, array_size,loadLevel);
  }
  
  // check if kernel execution generated and error
  CUT_CHECK_ERROR("Kernel execution failed");
  CUT_SAFE_CALL( cutStopTimer( timer));

  printf( "Processing time: %f (ms)\n", cutGetTimerValue(timer) );
  
  CUT_SAFE_CALL( cutDeleteTimer( timer));
  
  // allocate mem for the result on host side
  float* h_c = (float*) malloc( mem_size);
  // copy result from device to host
  CUDA_SAFE_CALL( cudaMemcpy( h_c, d_c, mem_size, cudaMemcpyDeviceToHost) );
  
  
  // compute reference solution
  float* ref_c = (float*) malloc( mem_size);
  for(int i=0; i<max_itr; i++)
    kernelHost( h_a, h_b, ref_c, array_size, loadLevel);
  
  // custom output handling when no regression test running
  // in this case check if the result is equivalent to the expected soluion
  CUTBoolean res = cutComparef( ref_c, h_c, array_size);
  printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED");
 
  // cleanup memory
  free( h_a);
  free( h_b);
  free( h_c);
  free( ref_c);
  CUDA_SAFE_CALL(cudaFree(d_a));
  CUDA_SAFE_CALL(cudaFree(d_b));
  CUDA_SAFE_CALL(cudaFree(d_c));
}