//
// include files
//

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

#include <hip/hip_runtime.h>

#include "helper_hip.h"

//
// kernel routine
// 

__global__ void my_first_kernel(float *x) {
  int tid = threadIdx.x + blockDim.x*blockIdx.x;

  x[tid] = (float) threadIdx.x;
}

//
// main code
//

int main(int argc, const char **argv) {
  float *h_x, *d_x;
  int   nblocks, nthreads, nsize, n;

  // initialise card

  findHipDevice(argc, argv);

  // set number of blocks, and threads per block

  nblocks  = 2;
  nthreads = 8;
  nsize    = nblocks*nthreads;

  // allocate memory for array

  h_x = (float *) malloc(nsize*sizeof(float));
  checkHipErrors( hipMalloc((void **) &d_x, nsize*sizeof(float)) );

  // execute kernel
  
  my_first_kernel<<<nblocks,nthreads>>>(d_x);
  getLastHipError("my_first_kernel execution failed\n");

  // copy back results and print them out

  checkHipErrors( hipMemcpy(h_x,d_x,nsize*sizeof(float),
                            hipMemcpyDefault) );

  for (n=0; n<nsize; n++) printf(" n,  x  =  %d  %f \n",n,h_x[n]);

  // free memory 

  checkHipErrors( hipFree(d_x) );
  free(h_x);

  // HIP exit -- needed to flush printf write buffer

  checkHipErrors( hipDeviceReset() );

  return 0;
}