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

/* THIS IS OPEN SOURCE CODE */

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

/**

 * @file    HelloWorld.c

 * CVS:     $Id$

 * @author Asim YarKhan (yarkhan@icl.utk.edu) HelloWorld altered to test power capping (October 2017) 

 * @author Heike Jagode (jagode@icl.utk.edu)

 * Mods:  <your name here> <your email address>

 *

 * @brief

 * This file is a very simple HelloWorld C example which serves

 * (together with its Makefile) as a guideline on how to add tests to

 * components.  This file tests the ability to do power control using

 * NVML.

 * The papi configure and papi Makefile will take care of the

 * compilation of the component tests (if all tests are added to a

 * directory named 'tests' in the specific component dir).  See

 * components/README for more details.

 *

 * The string "Hello World!" is mangled and then restored.

 */

#include <cuda.h>

#include <stdio.h>

#include "papi.h"

#include "papi_test.h"

#define PAPI

// Prototypes

__global__ void helloWorld( char* );

// Host function

int main( int argc, char** argv )

{

#ifdef PAPI

#define NUM_EVENTS 1

    int retval, i;

    int EventSet = PAPI_NULL;

    long long values[NUM_EVENTS];

    /* REPLACE THE EVENT NAME 'PAPI_FP_OPS' WITH A CUDA EVENT

       FOR THE CUDA DEVICE YOU ARE RUNNING ON.

       RUN papi_native_avail to get a list of CUDA events that are

       supported on your machine */

    // e.g. on a P100 nvml:::Tesla_P100-SXM2-16GB:power

    char *EventName[NUM_EVENTS];

    int events[NUM_EVENTS];

    int eventCount = 0;

    const PAPI_component_info_t *cmpinfo;

    char event_name[PAPI_MAX_STR_LEN];

    /* PAPI Initialization */

    retval = PAPI_library_init( PAPI_VER_CURRENT );

    if( retval != PAPI_VER_CURRENT ) fprintf( stderr, "PAPI_library_init failed\n" );

    printf( "PAPI_VERSION : %4d %6d %7d\n",

            PAPI_VERSION_MAJOR( PAPI_VERSION ),

            PAPI_VERSION_MINOR( PAPI_VERSION ),

            PAPI_VERSION_REVISION( PAPI_VERSION ) );

    int numcmp = PAPI_num_components();

    // printf( "Searching for nvml component among %d components\n", numcmp );

    int cid = 0;

    for( cid=0; cid

        cmpinfo = PAPI_get_component_info( cid );

        // printf( "Component %d (%d): %s: %d events\n", cid, cmpinfo->CmpIdx, cmpinfo->name, cmpinfo->num_native_events );

        if ( cmpinfo == NULL )

            test_fail( __FILE__, __LINE__,"PAPI_get_component_info failed\n",-1 );

        else if ( strstr( cmpinfo->name, "nvml" ) )

            break;

    }

    if ( cid==numcmp )

        test_skip( __FILE__, __LINE__,"Component nvml is not present\n",-1 );

    printf( "nvml component found: Component Index %d: %s: %d events\n", cmpinfo->CmpIdx, cmpinfo->name, cmpinfo->num_native_events );

    if ( cmpinfo->disabled )

        test_skip( __FILE__,__LINE__,"Component nvml is disabled", 0 );

    int code = PAPI_NATIVE_MASK;

    int ii=0;

    int event_modifier = PAPI_ENUM_FIRST;

    for ( ii=0; ii<cmpinfo->num_native_events; ii++ ) {

        retval = PAPI_enum_cmp_event( &code, event_modifier, cid );

        event_modifier = PAPI_ENUM_EVENTS;

        if ( retval != PAPI_OK ) test_fail( __FILE__, __LINE__, "PAPI_event_code_to_name", retval );

        retval = PAPI_event_code_to_name( code, event_name );

        // printf( "Look at event %d %d %s \n", ii, code, event_name );

        if ( strstr( event_name, "power_management_limit" ) )

            break;

    }

    if ( ii==cmpinfo->num_native_events )

        test_skip( __FILE__,__LINE__,"Component nvml does not have a power_management_limit event", 0 );

    printf( "nvml power_management_limit event found (%s)\n", event_name );

    EventName[0] = event_name;

    /* convert PAPI native events to PAPI code */

    for( i = 0; i < NUM_EVENTS; i++ ) {

        retval = PAPI_event_name_to_code( ( char * )EventName[i], &events[i] );

        if( retval != PAPI_OK )

            test_fail( __FILE__,__LINE__,"PAPI_event_name_to_code failed", retval );

        eventCount++;

        // printf( "Event: %s: Code: %#x\n", EventName[i], events[i] );

    }

    /* if we did not find any valid events, just report test failed. */

    if ( eventCount == 0 )

        test_skip( __FILE__,__LINE__,"No valid events found", retval );

    retval = PAPI_create_eventset( &EventSet );

    if( retval != PAPI_OK )

        test_fail( __FILE__,__LINE__,"PAPI_create_eventset failed", retval );

    retval = PAPI_add_events( EventSet, events, eventCount );

    if( retval != PAPI_OK )

        test_fail( __FILE__,__LINE__,"PAPI_add_events failed", retval );

#endif

    int j;

    int device_count;

    int cuda_device;

    cudaGetDeviceCount( &device_count );

    printf( "Found %d cuda devices\n", device_count );

///////////////////////   AYK

    for ( cuda_device = 0; cuda_device < device_count; cuda_device++ ) {

        // for ( cuda_device = 0; cuda_device < 1; cuda_device++ ) {

        printf( "cuda_device %d is being used\n", cuda_device );

        cudaSetDevice( cuda_device );

#ifdef PAPI

        retval = PAPI_start( EventSet );

        if( retval != PAPI_OK )

            test_fail( __FILE__,__LINE__,"PAPI_start failed", retval );

        retval = PAPI_read( EventSet, values );

        if( retval != PAPI_OK ) fprintf( stderr, "PAPI_read failed\n" );

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

            printf( "%s = %lld (read initial power management limit)\n", EventName[i], values[i]);

        long long int initial_power_management_limit = values[0];

        if ( cuda_device==0 ) {

            printf("On device_num %d the power_management_limit is going to be reduced by 30\n", cuda_device);

            // values[0] = 235000

            values[0] = initial_power_management_limit - 30;

            retval = PAPI_write( EventSet, values );

            if ( retval!=PAPI_OK ) {

                test_skip( __FILE__,__LINE__,"Attempted write of power_management_limit failed:  Possible reasons: Insufficient permissions; Power management unavailable. Outside min/max limits", retval );

            } else {

                printf( "Set power_management_limit to %llu milliWatts\n", values[0] );

            }

        }

#endif

        // desired output

        char str[] = "Hello World!";

        // mangle contents of output

        // the null character is left intact for simplicity

        for(j = 0; j < 12; j++) 

            str[j] -= j;

        printf( "This mangled string need to be fixed=%s\n", str );

        // allocate memory on the device

        char *d_str;

        size_t size = sizeof( str );

        cudaMalloc( ( void** )&d_str, size );

        // copy the string to the device

        cudaMemcpy( d_str, str, size, cudaMemcpyHostToDevice );

        // set the grid and block sizes

        dim3 dimGrid( 2 ); // one block per word

        dim3 dimBlock( sizeof( str )/2 ); // one thread per character

        // invoke the kernel

        helloWorld<<< dimGrid, dimBlock >>>( d_str );

        // retrieve the results from the device

        cudaMemcpy( str, d_str, size, cudaMemcpyDeviceToHost );

        // free up the allocated memory on the device

        cudaFree( d_str );

        printf( "Device %d Unmangled string = %s\n", cuda_device, str );

#ifdef PAPI

        if ( cuda_device==0 ) {

            retval = PAPI_read( EventSet, values );

            if( retval != PAPI_OK ) fprintf( stderr, "PAPI_read failed\n" );

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

                printf( "%s = %lld (read power management limit after reducing it... was it reduced?) \n", EventName[i], values[i] );

            if ( values[0] != initial_power_management_limit - 30 ) {

                printf( "Mismatch: power_management_limit on device %d set to %llu but read as %llu\n", cuda_device, initial_power_management_limit-30, values[0] );

                test_fail( __FILE__,__LINE__,"Mismatch: power_management_limit on device set to one value but read as a different value", -1 );

            }

            // AYK papi_reset

            long long resetvalues[NUM_EVENTS];

            resetvalues[0] = initial_power_management_limit;

            retval = PAPI_write( EventSet, resetvalues );

            retval = PAPI_stop( EventSet, values );

        }

#endif

    }

    test_pass( __FILE__);

    return 0;

}

// Device kernel

__global__ void

helloWorld( char* str )

{

    // determine where in the thread grid we are

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // unmangle output

    str[idx] += idx;

}