/* 

 * Copyright 2015-2016 NVIDIA Corporation. All rights reserved.

 *

 * Sample to demonstrate use of NVlink CUPTI APIs

 */

#include <stdio.h>

#include <stdlib.h>

#include <string.h>

#include <cuda.h>

#include <cupti.h>

#ifdef PAPI

#include "papi.h"

#endif

#define CUPTI_CALL(call)                                                \

    do {                                                                \

        CUptiResult _status = call;                                     \

        if (_status != CUPTI_SUCCESS) {                                 \

            const char *errstr;                                         \

            cuptiGetResultString(_status, &errstr);                     \

            fprintf(stderr, "%s:%d: error: function %s failed with error %s.\n", \

                    __FILE__, __LINE__, #call, errstr);                 \

            exit(-1);                                                   \

        }                                                               \

    } while (0)

#define DRIVER_API_CALL(apiFuncCall)                                    \

    do {                                                                \

        CUresult _status = apiFuncCall;                                 \

        if (_status != CUDA_SUCCESS) {                                  \

            fprintf(stderr, "%s:%d: error: function %s failed with error %d.\n", \

                    __FILE__, __LINE__, #apiFuncCall, _status);         \

            exit(-1);                                                   \

        }                                                               \

    } while (0)

#define RUNTIME_API_CALL(apiFuncCall)                                   \

    do {                                                                \

        cudaError_t _status = apiFuncCall;                              \

        if (_status != cudaSuccess) {                                   \

            fprintf(stderr, "%s:%d: error: function %s failed with error %s.\n", \

                    __FILE__, __LINE__, #apiFuncCall, cudaGetErrorString(_status)); \

            exit(-1);                                                   \

        }                                                               \

    } while (0)

#define MEMORY_ALLOCATION_CALL(var)                                     \

    do {                                                                \

        if (var == NULL) {                                              \

            fprintf(stderr, "%s:%d: Error: Memory Allocation Failed \n", \

                    __FILE__, __LINE__);                                \

            exit(-1);                                                   \

        }                                                               \

    } while (0)

#define MAX_DEVICES    (32)

#define BLOCK_SIZE     (1024)

#define GRID_SIZE      (512)

#define BUF_SIZE       (32 * 1024)

#define ALIGN_SIZE     (8)

#define SUCCESS        (0)

#define NUM_METRIC     (4)

#define NUM_EVENTS     (2)

#define MAX_SIZE       (64*1024*1024)   // 64 MB

#define NUM_STREAMS    (6)      // gp100 has 6 physical copy engines

CUpti_ActivityNvLink *nvlinkRec = NULL;

int cpuToGpu = 0;

int gpuToGpu = 0;

int cpuToGpuAccess = 0;

int gpuToGpuAccess = 0;

extern "C" __global__ void test_nvlink_bandwidth(float *src, float *dst)

{

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

    dst[idx] = src[idx] * 2.0f;

}

static void printActivity(CUpti_Activity * record)

{

    if(record->kind == CUPTI_ACTIVITY_KIND_NVLINK) {

        nvlinkRec = (CUpti_ActivityNvLink *) record;

        // printf("typeDev0 %d, typeDev1 %d, sysmem %d, peer %d, physical links %d, portdev0 %d, %d, %d, %d, portDev1 %d, %d, %d, %d, bandwidth %llu\n", nvlinkRec->typeDev0, nvlinkRec->typeDev1, ((nvlinkRec->flag & CUPTI_LINK_FLAG_SYSMEM_ACCESS) ? 1 : 0), ((nvlinkRec->flag & CUPTI_LINK_FLAG_PEER_ACCESS) ? 1 : 0), nvlinkRec->physicalNvLinkCount, nvlinkRec->portDev0[0], nvlinkRec->portDev0[1], nvlinkRec->portDev0[2], nvlinkRec->portDev0[3], nvlinkRec->portDev1[0], nvlinkRec->portDev1[1], nvlinkRec->portDev1[2], nvlinkRec->portDev1[3], (long long unsigned int) nvlinkRec->bandwidth);

        cpuToGpuAccess |= (nvlinkRec->flag & CUPTI_LINK_FLAG_SYSMEM_ACCESS);

        gpuToGpuAccess |= (nvlinkRec->flag & CUPTI_LINK_FLAG_PEER_ACCESS);

    } else {

        printf("Error : Unexpected CUPTI activity kind.\nExpected Activity kind : CUPTI_ACTIVITY_KIND_NVLINK\n");

    }

}

static void CUPTIAPI bufferRequested(uint8_t ** buffer, size_t * size, size_t * maxNumRecords)

{

    *size = BUF_SIZE + ALIGN_SIZE;

    *buffer = (uint8_t *) calloc(1, *size);

    MEMORY_ALLOCATION_CALL(*buffer);

    *maxNumRecords = 0;

}

static void CUPTIAPI bufferCompleted(CUcontext ctx, uint32_t streamId, uint8_t * buffer, size_t size, size_t validSize)

{

    CUptiResult status;

    CUpti_Activity *record = NULL;

    do {

        status = cuptiActivityGetNextRecord(buffer, validSize, &record);

        if(status == CUPTI_SUCCESS) {

            printActivity(record);

        } else if(status == CUPTI_ERROR_MAX_LIMIT_REACHED) {

            break;

        } else {

            CUPTI_CALL(status);

        }

    } while(1);

    size_t dropped;

    CUPTI_CALL(cuptiActivityGetNumDroppedRecords(ctx, streamId, &dropped));

    if(dropped != 0) {

        printf("Dropped %u activity records\n", (unsigned int) dropped);

    }

}

#define DIM(x) (sizeof(x)/sizeof(*(x)))

void calculateSize(char *result, uint64_t size)

{

    int i;

    const char *sizes[] = { "TB", "GB", "MB", "KB", "B" };

    uint64_t exbibytes = 1024ULL * 1024ULL * 1024ULL * 1024ULL;

    uint64_t multiplier = exbibytes;

    for(i = 0; (unsigned) i < DIM(sizes); i++, multiplier /= (uint64_t) 1024) {

        if(size < multiplier)

            continue;

        sprintf(result, "%.1f %s", (float) size / multiplier, sizes[i]);

        return;

    }

    strcpy(result, "0");

    return;

}

void readMetricValue(CUpti_EventGroup eventGroup, uint32_t numEvents, CUdevice dev, CUpti_MetricID * metricId, uint64_t timeDuration, CUpti_MetricValue * metricValue)

{

    size_t bufferSizeBytes, numCountersRead;

    uint64_t *eventValueArray = NULL;

    CUpti_EventID *eventIdArray;

    size_t arraySizeBytes = 0;

    size_t numTotalInstancesSize = 0;

    uint64_t numTotalInstances = 0;

    uint64_t *aggrEventValueArray = NULL;

    size_t aggrEventValueArraySize;

    uint32_t i = 0, j = 0;

    CUpti_EventDomainID domainId;

    size_t domainSize;

    domainSize = sizeof(CUpti_EventDomainID);

    CUPTI_CALL(cuptiEventGroupGetAttribute(eventGroup, CUPTI_EVENT_GROUP_ATTR_EVENT_DOMAIN_ID, &domainSize, (void *) &domainId));

    numTotalInstancesSize = sizeof(uint64_t);

    CUPTI_CALL(cuptiDeviceGetEventDomainAttribute(dev, domainId, CUPTI_EVENT_DOMAIN_ATTR_TOTAL_INSTANCE_COUNT, &numTotalInstancesSize, (void *) &numTotalInstances));

    arraySizeBytes = sizeof(CUpti_EventID) * numEvents;

    bufferSizeBytes = sizeof(uint64_t) * numEvents * numTotalInstances;

    eventValueArray = (uint64_t *) malloc(bufferSizeBytes);

    MEMORY_ALLOCATION_CALL(eventValueArray);

    eventIdArray = (CUpti_EventID *) malloc(arraySizeBytes);

    MEMORY_ALLOCATION_CALL(eventIdArray);

    aggrEventValueArray = (uint64_t *) calloc(numEvents, sizeof(uint64_t));

    MEMORY_ALLOCATION_CALL(aggrEventValueArray);

    aggrEventValueArraySize = sizeof(uint64_t) * numEvents;

    CUPTI_CALL(cuptiEventGroupReadAllEvents(eventGroup, CUPTI_EVENT_READ_FLAG_NONE, &bufferSizeBytes, eventValueArray, &arraySizeBytes, eventIdArray, &numCountersRead));

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

        for(j = 0; j < numTotalInstances; j++) {

            aggrEventValueArray[i] += eventValueArray[i + numEvents * j];

            printf("For event %d instance %d value %ul aggregate %d = %ul\n", i, j,  eventValueArray[i + numEvents * j], i, aggrEventValueArray[i]);

        }

    }

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

        CUPTI_CALL(cuptiMetricGetValue(dev, metricId[i], arraySizeBytes, eventIdArray, aggrEventValueArraySize, aggrEventValueArray, timeDuration, &metricValue[i]));

    }

    free(eventValueArray);

    free(eventIdArray);

}

// Print metric value, we format based on the value kind

int printMetricValue(CUpti_MetricID metricId, CUpti_MetricValue metricValue, const char *metricName)

{

    CUpti_MetricValueKind valueKind;

    char str[64];

    size_t valueKindSize = sizeof(valueKind);

    CUPTI_CALL(cuptiMetricGetAttribute(metricId, CUPTI_METRIC_ATTR_VALUE_KIND, &valueKindSize, &valueKind));

    switch (valueKind) {

    case CUPTI_METRIC_VALUE_KIND_DOUBLE:

        printf("%s = %f ", metricName, metricValue.metricValueDouble);

        calculateSize(str, (uint64_t) metricValue.metricValueDouble);

        printf("%s\n", str);

        break;

    case CUPTI_METRIC_VALUE_KIND_UINT64:

        printf("%s = %lu ", metricName, metricValue.metricValueUint64);

        calculateSize(str, (uint64_t) metricValue.metricValueUint64);

        printf("%s\n", str);

        break;

    case CUPTI_METRIC_VALUE_KIND_INT64:

        printf("%s = %ld ", metricName, metricValue.metricValueInt64);

        calculateSize(str, (uint64_t) metricValue.metricValueInt64);

        printf("%s\n", str);

        break;

    case CUPTI_METRIC_VALUE_KIND_THROUGHPUT:

        printf("%s = %f ", metricName, metricValue.metricValueThroughput);

        calculateSize(str, (uint64_t) metricValue.metricValueThroughput);

        printf("%s/Sec\n", str);

        break;

    default:

        fprintf(stderr, "error: unknown value kind\n");

        return -1;

    }

    return 0;

}

void testCpuToGpu(CUpti_EventGroup * eventGroup, CUdeviceptr * pDevBuffer, float **pHostBuffer, size_t bufferSize, cudaStream_t * cudaStreams, uint64_t * timeDuration, int numEventGroup)

{

    int i;

#ifdef CUPTI_ONLY

    uint32_t value = 1;

    uint64_t startTimestamp, endTimestamp;

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

        CUPTI_CALL(cuptiEventGroupEnable(eventGroup[i]));

        CUPTI_CALL(cuptiEventGroupSetAttribute(eventGroup[i], CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES, sizeof(uint32_t), (void *) &value));

    }

    CUPTI_CALL(cuptiGetTimestamp(&startTimestamp));

#endif                          // CUPTI_ONLY

    // Unidirectional copy H2D

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

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer[i], pHostBuffer[i], bufferSize, cudaMemcpyHostToDevice, cudaStreams[i]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

    // Unidirectional copy D2H

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

        RUNTIME_API_CALL(cudaMemcpyAsync(pHostBuffer[i], (void *) pDevBuffer[i], bufferSize, cudaMemcpyDeviceToHost, cudaStreams[i]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

    // Bidirectional copy

    for(i = 0; i < NUM_STREAMS; i += 2) {

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer[i], pHostBuffer[i], bufferSize, cudaMemcpyHostToDevice, cudaStreams[i]));

        RUNTIME_API_CALL(cudaMemcpyAsync(pHostBuffer[i + 1], (void *) pDevBuffer[i + 1], bufferSize, cudaMemcpyDeviceToHost, cudaStreams[i + 1]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

#ifdef CUPTI_ONLY

    CUPTI_CALL(cuptiGetTimestamp(&endTimestamp));

    *timeDuration = endTimestamp - startTimestamp;

#endif                          // CUPTI_ONLY

}

void testGpuToGpu_part1(CUpti_EventGroup * eventGroup, CUdeviceptr * pDevBuffer0, CUdeviceptr * pDevBuffer1, float **pHostBuffer, size_t bufferSize, cudaStream_t * cudaStreams, uint64_t * timeDuration, int numEventGroup)

{

    int i;

#ifdef CUPTI_ONLY

    uint32_t value = 1;

    uint64_t startTimestamp, endTimestamp;

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

        CUPTI_CALL(cuptiEventGroupEnable(eventGroup[i]));

        CUPTI_CALL(cuptiEventGroupSetAttribute(eventGroup[i], CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES, sizeof(uint32_t), (void *) &value));

    }

#endif                          // CUPTI_ONLY

    RUNTIME_API_CALL(cudaSetDevice(0));

    RUNTIME_API_CALL(cudaDeviceEnablePeerAccess(1, 0));

    RUNTIME_API_CALL(cudaSetDevice(1));

    RUNTIME_API_CALL(cudaDeviceEnablePeerAccess(0, 0));

    // Unidirectional copy H2D

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

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer0[i], pHostBuffer[i], bufferSize, cudaMemcpyHostToDevice, cudaStreams[i]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

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

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer1[i], pHostBuffer[i], bufferSize, cudaMemcpyHostToDevice, cudaStreams[i]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

}

void testGpuToGpu_part2(CUpti_EventGroup * eventGroup, CUdeviceptr * pDevBuffer0, CUdeviceptr * pDevBuffer1, float **pHostBuffer, size_t bufferSize, cudaStream_t * cudaStreams, uint64_t * timeDuration, int numEventGroup)

{

    int i;

#ifdef CUPTI_ONLY

    uint32_t value = 1;

    uint64_t startTimestamp, endTimestamp;

    CUPTI_CALL(cuptiGetTimestamp(&startTimestamp));

#endif

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

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer0[i], (void *) pDevBuffer1[i], bufferSize, cudaMemcpyDeviceToDevice, cudaStreams[i]));

        //printf("Copy %zu stream %d to devBuffer0 from devBuffer1 \n", bufferSize, i);

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

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

        RUNTIME_API_CALL(cudaMemcpyAsync((void *) pDevBuffer1[i], (void *) pDevBuffer0[i], bufferSize, cudaMemcpyDeviceToDevice, cudaStreams[i]));

        // printf("Copy %zu stream %d to devBuffer0 from devBuffer1 \n", bufferSize, i);

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

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

        test_nvlink_bandwidth <<< GRID_SIZE, BLOCK_SIZE >>> ((float *) pDevBuffer1[i], (float *) pDevBuffer0[i]);

        // printf("test_nvlink_bandwidth stream %d \n", i);

    }

#ifdef CUPTI_ONLY

    CUPTI_CALL(cuptiGetTimestamp(&endTimestamp));

    *timeDuration = endTimestamp - startTimestamp;

#endif                          // CUPTI_ONLY

}

static void printUsage()

{

    printf("usage: Demonstrate use of NVlink CUPTI APIs\n");

    printf("       -help           : display help message\n");

    printf("       --cpu-to-gpu    : Show results for data transfer between CPU and GPU \n");

    printf("       --gpu-to-gpu    : Show results for data transfer between two GPUs \n");

}

void parseCommandLineArgs(int argc, char *argv[])

{

    if(argc != 2) {

        printf("Invalid number of options\n");

        exit(0);

    }

    if(strcmp(argv[1], "--cpu-to-gpu") == 0) {

        cpuToGpu = 1;

    } else if(strcmp(argv[1], "--gpu-to-gpu") == 0) {

        gpuToGpu = 1;

    } else if((strcmp(argv[1], "--help") == 0) || (strcmp(argv[1], "-help") == 0) || (strcmp(argv[1], "-h") == 0)) {

        printUsage();

        exit(0);

    } else {

        cpuToGpu = 1;

    }

}

int main(int argc, char *argv[])

{

    int deviceCount = 0, i = 0, numEventGroup = 0;

    size_t bufferSize = 0, freeMemory = 0, totalMemory = 0;

    CUcontext ctx;

    char str[64];

    CUdeviceptr pDevBuffer0[NUM_STREAMS];

    CUdeviceptr pDevBuffer1[NUM_STREAMS];

    float *pHostBuffer[NUM_STREAMS];

    cudaStream_t cudaStreams[NUM_STREAMS] = { 0 };

    cudaDeviceProp prop[MAX_DEVICES];

    uint64_t timeDuration;

    CUpti_EventGroup eventGroup[32];

#ifdef CUPTI_ONLY

    CUpti_MetricID metricId[NUM_METRIC];

    uint32_t numEvents[NUM_METRIC];

    CUpti_MetricValue metricValue[NUM_METRIC];

    // Adding nvlink Metrics.

    const char *metricName[NUM_METRIC] = { 

        "nvlink_total_data_transmitted",

        "nvlink_total_data_received",

        "nvlink_transmit_throughput",

        "nvlink_receive_throughput"

    };

#endif                          // CUPTI_ONLY

    // Parse command line arguments

    parseCommandLineArgs(argc, argv);

    CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_NVLINK));

    CUPTI_CALL(cuptiActivityRegisterCallbacks(bufferRequested, bufferCompleted));

    DRIVER_API_CALL(cuInit(0));

    RUNTIME_API_CALL(cudaGetDeviceCount(&deviceCount));

    printf("There are %d devices.\n", deviceCount);

    if(deviceCount == 0) {

        printf("There is no device supporting CUDA.\n");

        exit(-1);

    }

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

        RUNTIME_API_CALL(cudaGetDeviceProperties(&prop[i], i));

        printf("CUDA Device %d Name: %s\n", i, prop[i].name);

    }

    // Set memcpy size based on available device memory

    RUNTIME_API_CALL(cudaMemGetInfo(&freeMemory, &totalMemory));

    printf("Total Device Memory available : ");

    calculateSize(str, (uint64_t) totalMemory);

    printf("%s\n", str);

    bufferSize = MAX_SIZE < (freeMemory / 4) ? MAX_SIZE : (freeMemory / 4);

    bufferSize = bufferSize/2;

    printf("Memcpy size is set to %llu B (%llu MB)\n", (unsigned long long) bufferSize, (unsigned long long) bufferSize / (1024 * 1024));

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

        RUNTIME_API_CALL(cudaStreamCreate(&cudaStreams[i]));

    }

    RUNTIME_API_CALL(cudaDeviceSynchronize());

    // Nvlink-topology Records are generated even before cudaMemcpy API is called.

    CUPTI_CALL(cuptiActivityFlushAll(0));

    // Transfer Data between Host And Device, if Nvlink is Present

    // Check condition : nvlinkRec->flag & CUPTI_LINK_FLAG_SYSMEM_ACCESS

    // True : Nvlink is present between CPU & GPU

    // False : Nvlink is not present.

    if((nvlinkRec) && (((cpuToGpu) && (cpuToGpuAccess)) || ((gpuToGpu) && (gpuToGpuAccess)))) {

#ifdef CUPTI_ONLY

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

            CUPTI_CALL(cuptiMetricGetIdFromName(0, metricName[i], &metricId[i]));

            CUPTI_CALL(cuptiMetricGetNumEvents(metricId[i], &numEvents[i]));

        }

#endif                          // CUPTI_ONLY

        DRIVER_API_CALL(cuCtxCreate(&ctx, 0, 0));

#ifdef PAPI

        printf("Setup PAPI counters internally (PAPI)\n");

        int EventSet = PAPI_NULL;

        long long values[MAX_DEVICES * NUM_METRIC];

        char *EventName[MAX_DEVICES * NUM_METRIC];

        int eventCount;

        int retval, ee;

        /* PAPI Initialization */

        retval = PAPI_library_init(PAPI_VER_CURRENT);

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

        printf("PAPI version: %d.%d.%d\n", PAPI_VERSION_MAJOR(PAPI_VERSION), PAPI_VERSION_MINOR(PAPI_VERSION), PAPI_VERSION_REVISION(PAPI_VERSION));

        retval = PAPI_create_eventset(&EventSet);

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

        const char *EventEndings[NUM_METRIC] = {

            "cuda:::metric:nvlink_total_data_transmitted",

            "cuda:::metric:nvlink_total_data_received",

            "cuda:::metric:nvlink_transmit_throughput",

            "cuda:::metric:nvlink_receive_throughput",

        };

        // Add events at a GPU specific level ... eg cuda:::metric:nvlink_total_data_transmitted:device=0

        // Just profile devices to match the CUPTI example

        char tmpEventName[1024];

        eventCount = 0;

        for(i = 0; i < 1; i++) { // only profile device 0

            printf("Set device to %d\n", i);

            for(ee = 0; ee < NUM_METRIC; ee++) {

                snprintf(tmpEventName, 1024, "%s:device=%d\0", EventEndings[ee], i);

                printf("Trying to add event %s to GPU %d in PAPI...", tmpEventName, i);

                retval = PAPI_add_named_event(EventSet, tmpEventName);

                if(retval == PAPI_OK) {

                    printf("Added event\n");

                    EventName[eventCount] = strdup(tmpEventName);

                    eventCount++;

                } else {

                    printf("Could not add event\n");

                }

            }

        }

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

            values[i] = -1;

#endif                          // PAPI_ONLY

#ifdef CUPTI_ONLY

        CUpti_EventGroupSets *passes = NULL;

        int j = 0;

        CUPTI_CALL(cuptiMetricCreateEventGroupSets(ctx, (sizeof metricId), metricId, &passes));

        // EventGroups required to profile Nvlink metrics.

        for(i = 0; i < (signed) passes->numSets; i++) {

            for(j = 0; j < (signed) passes->sets[i].numEventGroups; j++) {

                eventGroup[numEventGroup] = passes->sets[i].eventGroups[j];

                if(!eventGroup[numEventGroup]) {

                    printf("\n eventGroup initialization failed \n");

                    exit(-1);

                }

                numEventGroup++;

            }

        }

        CUPTI_CALL(cuptiSetEventCollectionMode(ctx, CUPTI_EVENT_COLLECTION_MODE_CONTINUOUS));

#endif                          // CUPTI_ONLY

        // ===== Allocate Memory =====================================

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

            RUNTIME_API_CALL(cudaMalloc((void **) &pDevBuffer0[i], bufferSize));

            pHostBuffer[i] = (float *) malloc(bufferSize);

            MEMORY_ALLOCATION_CALL(pHostBuffer[i]);

        }

        if(cpuToGpu) {

#ifdef PAPI

            retval = PAPI_start( EventSet );

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

#endif// PAPI

            testCpuToGpu(eventGroup, pDevBuffer0, pHostBuffer, bufferSize, cudaStreams, &timeDuration, numEventGroup);

#ifdef PAPI

            retval = PAPI_stop(EventSet, values);

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

#endif 

            printf("Data tranferred between CPU & Device%d : \n", (int) nvlinkRec->typeDev0);

        } else if(gpuToGpu) {

            RUNTIME_API_CALL(cudaSetDevice(1));

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

                RUNTIME_API_CALL(cudaMalloc((void **) &pDevBuffer1[i], bufferSize));

            testGpuToGpu_part1(eventGroup, pDevBuffer0, pDevBuffer1, pHostBuffer, bufferSize, cudaStreams, &timeDuration, numEventGroup);

#ifdef PAPI

            retval = PAPI_start( EventSet );

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

#endif

            testGpuToGpu_part2(eventGroup, pDevBuffer0, pDevBuffer1, pHostBuffer, bufferSize, cudaStreams, &timeDuration, numEventGroup);

#ifdef PAPI

            retval = PAPI_stop(EventSet, values);

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

#endif

            printf("Data tranferred between Device 0 & Device 1 : \n");

        }

#ifdef CUPTI_ONLY

        // Collect Nvlink Metric values for the data transfer via Nvlink for all the eventGroups.

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

            readMetricValue(eventGroup[i], NUM_EVENTS, 0, metricId, timeDuration, metricValue);

            CUPTI_CALL(cuptiEventGroupDisable(eventGroup[i]));

            CUPTI_CALL(cuptiEventGroupDestroy(eventGroup[i]));

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

                if(printMetricValue(metricId[i], metricValue[i], metricName[i]) != 0) {

                    printf("\n printMetricValue failed \n");

                    exit(-1);

                }

            }

        }

#endif                          // CUPTI_ONLY

#ifdef PAPI

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

            char str[64];                

            calculateSize(str, (uint64_t) values[i] );

            printf("PAPI %s %s \n", EventName[i], str);

        }

	retval = PAPI_cleanup_eventset(EventSet);

	if( retval != PAPI_OK )

            fprintf(stderr, "PAPI_cleanup_eventset failed\n");

	retval = PAPI_destroy_eventset(&EventSet);

	if (retval != PAPI_OK)

            fprintf(stderr, "PAPI_destroy_eventset failed\n");

	PAPI_shutdown();

#endif

    } else {

        printf("No Nvlink supported device found\n");

    }

    printf("Exit\n");

    return 0;

}