/*
 * Copyright 1993-2017 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.
 */

// Simple interop app demonstrating EGLImage + EGLSync interop with CUDA.
// Using EGLSync - CUDA Event interop one can achieve synchronization on GPU itself for GL-EGL-CUDA operations 
// instead of blocking CPU for synchronization.
// This app requires GLES 3.2 or higher

//---------------------------INCLUDES---------------------------------//
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <time.h>
#include "graphics_interface.h"
#include <cuda.h>
#include <helper_cuda_drvapi.h>
#include <cudaEGL.h>
#include <EGL/egl.h>
#include <EGL/eglext.h>
#include <GLES3/gl32.h>
#include "egl_common.h"

//---------------------------DEFINES---------------------------------//
#define MAX_ITR             100

#define FAILURE             0
#define SUCCESS             1
#define WAIVED              2

#define BLOCK_SIZE          16

#define GL_READ             0
#define GL_WRITE            1
//---------------------------MACROS---------------------------------//

// Error-checking wrapper around GL calls
#define GL_SAFE_CALL(call)              \
{                                   \
    GLenum err;                     \
    call;                           \
    err = glGetError();             \
    if(err != GL_NO_ERROR)          \
    {                               \
        fprintf(stderr, "%s:%d GL error: %d\n", __FILE__, __LINE__, err);    \
        cleanup(FAILURE);           \
    }                               \
}

#define GL_SAFE_CALL_NO_CLEANUP(call, err)  \
{                                       \
    GLenum status;                      \
    call;                               \
    status = glGetError();              \
    if(status != GL_NO_ERROR)           \
    {                                   \
        fprintf(stderr, "%s:%d GL error: %d\n", __FILE__, __LINE__, status);    \
        err = status;                   \
    }                                   \
}

// Error-checking wrapper around CUDA calls (taken from cutil.h)
#define CUDA_SAFE_CALL(call) do {                                           \
    cudaError err = call;                                                   \
    if( cudaSuccess != err) {                                               \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",       \
                __FILE__, __LINE__, cudaGetErrorString( err) );             \
        cleanup(FAILURE);                                                   \
    } } while (0)

#define CUDA_SAFE_CALL_NO_CLEANUP(call, err) do {                           \
    cudaError status = call;                                                \
    if( cudaSuccess != status) {                                            \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",       \
                __FILE__, __LINE__, cudaGetErrorString( status) );          \
        err = status;                                                       \
    } } while (0)


#if defined(EXTENSION_LIST)
EXTENSION_LIST(EXTLST_DECL)
    typedef void (*extlst_fnptr_t)(void);
    static struct {
        extlst_fnptr_t *fnptr;
        char const *name;
    } extensionList[] = { EXTENSION_LIST(EXTLST_ENTRY) };

int eglSetupExtensions(void)
{
    unsigned int i;

    for (i = 0; i < (sizeof(extensionList) / sizeof(*extensionList)); i++) 
    {
        *extensionList[i].fnptr = eglGetProcAddress(extensionList[i].name);
        if (*extensionList[i].fnptr == NULL) 
        {
            printf("Couldn't get address of %s()\n", extensionList[i].name);
            return 0;
        }
    }

    return 1;
}
#endif

#if defined(EXTENSION_LIST)
EXTENSION_LIST(EXTLST_EXTERN)
#endif

//------------------------GLOBAL VARIABLES--------------------------//

// GL texture
GLuint tex[2]    = {0};

// Used to catch unexpected termination from GLUT
int cleanExit = 0;

//Use CPU Sync or GPU sync; Default GPU
int useGpu = 1;

// CUDA Resource
CUgraphicsResource writeResource = NULL; 
CUgraphicsResource readResource = NULL; 
CUarray writeArray,readArray;
CUdevice device;
CUcontext context;

// Which device to run on
unsigned int dev =0;

//Default width, height, and iterations value 
int width  = 2048;
int height = 2048;
int itr    = MAX_ITR;

// Error check variable
__device__ static unsigned int numErrors = 0;

//-----------------------FUNCTION PROTOTYPES------------------------//

void            checkSync(int argc, char **argv);
int             parseCmdLine(int argc, char **argv);
void            printUsage(void);
void            cleanup(int status);
void            exitHandler(void);
void            printStatus(int status);
void            checkSyncOnCPU(void);
void            checkSyncOnGPU(EGLDisplay dpy);

__global__ void verify_and_update_kernel(CUsurfObject write, CUsurfObject read, char expected, char newval, int width, int height);
extern "C" cudaError_t cudaGetValueMismatch();

//-----------------------FUNCTION DEFINITIONS------------------------//

int main(int argc, char *argv[])
{
#if defined(__linux__)
    setenv ("DISPLAY", ":0", 0);
#endif

    parseCmdLine(argc, argv);
    atexit(exitHandler);

    checkSync(argc, argv);
    return 0;
}

int parseCmdLine(int argc, char **argv) 
{
    int i;
    for(i = 1; i < argc; i++) 
    {
        if(strcmp(argv[i], "-cpu") == 0) 
        {
            useGpu = 0;
        }

        if(strcmp(argv[i], "-h") == 0) 
        {
            printUsage();
            cleanup(SUCCESS);
        }

        if(strcmp(argv[i], "-width") == 0) 
        {
            ++i;
            if(i == argc) 
            {
                printf("width option must be followed by value\n");
                return FAILURE;
            }
            if(sscanf(argv[i], "%d", &width) != 1) 
            {
                printf("Error: invalid width value\n");
                return FAILURE;
            }
        }

        if(strcmp(argv[i], "-height") == 0) 
        {
            ++i;
            if(i == argc) 
            {
                printf("height option must be followed by value\n");
                return FAILURE;
            }
            if(sscanf(argv[i], "%d", &height) != 1) 
            {
                printf("Error: invalid height value\n");
                return FAILURE;
            }
        }
        if(strcmp(argv[i], "-itr") == 0) 
        {
            ++i;
            if(i == argc) 
            {
                printf("itr option must be followed by iteration value\n");
                return FAILURE;
            }
            if(sscanf(argv[i], "%d", &itr) != 1)
            {
                printf("Error: invalid iteration value\n");
                return FAILURE;
            }
        }
    }

    return SUCCESS;
}

void printUsage(void) 
{
    printf("Usage:\n");
    printf("\t-h\tPrint command line options\n");
    printf("\t-cpu\tSync on the CPU instead of the GPU\n");
    printf("\t-width w\tSet the width to w\n");
    printf("\t-height h\tSet the height to h\n");
    printf("\t-itr i\tSet number of iterations to i\n");
}

void checkSync(int argc, char **argv)
{
    int x, y;
    int bufferSize = width * height * 4;
    unsigned char *pSurf_read = NULL, *pSurf_write = NULL;

    CUresult status = CUDA_SUCCESS;

    // Init values for variables
    x = y = 0;

    if (CUDA_SUCCESS != (status = cuInit(0)))
    {
        printf("Failed to initialize CUDA\n");
    }
    device = findCudaDeviceDRV(argc, (const char **)argv);

    if (CUDA_SUCCESS !=  (status = cuCtxCreate(&context, 0, device)))
    {
        printf("failed to create CUDA context\n");
    }
    cuCtxPushCurrent(context);

#if (defined (__arm__) || defined (__aarch64__)) && defined (__linux__)
    graphics_setup_window(0, 0, width, height, "EGLSync_CUDA_Interop");
#endif

    pSurf_read = (unsigned char *)malloc(bufferSize);
    pSurf_write = (unsigned char *)malloc(bufferSize);
    if (pSurf_read == NULL || pSurf_write == NULL)
    {
        printf("malloc failed\n");
        cleanup(FAILURE);
    }

    for(x = 0; x < width; x++)
    {
        for(y = 0; y < height; y++)
        {
            pSurf_read[(y*width + x) * 4    ] = 1;
            pSurf_read[(y*width + x) * 4 + 1] = 1;
            pSurf_read[(y*width + x) * 4 + 2] = 1;
            pSurf_read[(y*width + x) * 4 + 3] = 1;
            pSurf_write[(y*width + x) * 4    ] = 0;
            pSurf_write[(y*width + x) * 4 + 1] = 0;
            pSurf_write[(y*width + x) * 4 + 2] = 0;
            pSurf_write[(y*width + x) * 4 + 3] = 0;
        }
    }

    // NOP call to error-check the above glut calls
    GL_SAFE_CALL({});

    //Init texture
    GL_SAFE_CALL(glGenTextures(2, tex));

    GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, tex[GL_READ]));
    GL_SAFE_CALL(glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST));
    GL_SAFE_CALL(glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST));
    GL_SAFE_CALL(glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, pSurf_read));
    GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, 0));


    GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, tex[GL_WRITE]));
    GL_SAFE_CALL(glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST));
    GL_SAFE_CALL(glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST));
    GL_SAFE_CALL(glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, pSurf_write));
    GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, 0));

    glFinish();

    EGLDisplay eglDisplayHandle = eglGetCurrentDisplay();
    if (eglDisplayHandle == EGL_NO_DISPLAY) {
        printf("eglDisplayHandle failed \n");
        cleanup(FAILURE);
    }
    else {
        printf("eglDisplay Handle created \n");
    }

    if(!eglSetupExtensions()) {
        printf("SetupExtentions failed \n");
        cleanup(FAILURE);
    }

    EGLContext eglCtx = eglGetCurrentContext();
    if (eglCtx == EGL_NO_CONTEXT) {
        printf("Context1 create failed with error %d\n", eglGetError());
        cleanup(FAILURE);
    }

    // Create the EGL_Image
    EGLint eglImgAttrs[] = { EGL_IMAGE_PRESERVED_KHR, EGL_TRUE, EGL_NONE, EGL_NONE };

    EGLImageKHR eglImage1 = eglCreateImageKHR(eglDisplayHandle, eglCtx, EGL_GL_TEXTURE_2D_KHR, (EGLClientBuffer)(intptr_t)tex[GL_READ], eglImgAttrs);
    if (eglImage1 == EGL_NO_IMAGE_KHR) {
        printf("EGLImage create failed for read texture with error %d\n", eglGetError());
        cleanup(FAILURE);
    }
    else {
        printf("EGLImage1 created \n");
    }

    EGLImageKHR eglImage2 = eglCreateImageKHR(eglDisplayHandle, eglCtx, EGL_GL_TEXTURE_2D_KHR, (EGLClientBuffer)(intptr_t)tex[GL_WRITE], eglImgAttrs);
    if (eglImage2 == EGL_NO_IMAGE_KHR) {
        printf("EGLImage create failed for write texture with error %d\n", eglGetError());
        cleanup(FAILURE);
    }
    else {
        printf("EGLImage2 created \n");
    }

    glFinish();

    status = cuGraphicsEGLRegisterImage(&writeResource, eglImage1, CU_GRAPHICS_REGISTER_FLAGS_NONE);
    if (status != CUDA_SUCCESS) {
        printf("cuGraphicsEGLRegisterImage failed with Texture 1\n");
        cleanup(FAILURE);
    }
    else {
        printf("cuGraphicsEGLRegisterImage Passed, writeResource created with texture 1\n");
    }

    status = cuGraphicsSubResourceGetMappedArray(&writeArray, writeResource, 0, 0);
    if (status != CUDA_SUCCESS) {
        printf ("cuGraphicsSubResourceGetMappedArray failed for writeResource with texture 1\n");
        cleanup(FAILURE);
    }

    status = cuGraphicsEGLRegisterImage(&readResource, eglImage2, CU_GRAPHICS_REGISTER_FLAGS_NONE);
    if (status != CUDA_SUCCESS) {
        printf("cuGraphicsEGLRegisterImage failed for readResource with Texture 2\n");
        cleanup(FAILURE);
    }
    else {
        printf("cuGraphicsEGLRegisterImage Passed, readResource created with texture 2\n");
    }

    status = cuGraphicsSubResourceGetMappedArray(&readArray, readResource, 0, 0);
    if (status != CUDA_SUCCESS) {
        printf ("cuGraphicsSubResourceGetMappedArray failed for texture 2\n");
        cleanup(FAILURE);
    }

    if(useGpu) {
        printf("Using GPU Sync path\n");
        checkSyncOnGPU(eglDisplayHandle);
    }
    else {
        printf("Using CPU Sync path\n");
        checkSyncOnCPU();
    }

    free(pSurf_read);
    free(pSurf_write);
    cleanup(SUCCESS);
}

void checkSyncOnCPU(void) 
{
    int z = 0;
    unsigned char expectedData, newData;
    CUresult status = CUDA_SUCCESS;
    CUDA_RESOURCE_DESC wdsc,rdsc;
    memset(&wdsc, 0, sizeof(wdsc));
    memset(&rdsc, 0, sizeof(rdsc));

    expectedData = 0;
    newData = 1;

    wdsc.resType = CU_RESOURCE_TYPE_ARRAY;
    wdsc.res.array.hArray = writeArray;
    CUsurfObject writeSurface;
    rdsc.resType = CU_RESOURCE_TYPE_ARRAY;
    rdsc.res.array.hArray = readArray;
    CUsurfObject readSurface;

    status = cuSurfObjectCreate(&writeSurface, &wdsc);
    if (status != CUDA_SUCCESS)
    {   
        printf("Surface bounding failed with status %d\n",status);
        cleanup(FAILURE);
    }
    status = cuSurfObjectCreate(&readSurface, &rdsc);
    if (status != CUDA_SUCCESS)
    {
        printf("Surface bounding failed\n");
        cleanup(FAILURE);
    }

    for (z = 0; z < itr; z++)
    {
        //GL call to copy from read texture to write texture
        GL_SAFE_CALL(glCopyImageSubData(tex[GL_READ], GL_TEXTURE_2D, 0, 0, 0, 0,
                    tex[GL_WRITE], GL_TEXTURE_2D, 0, 0, 0, 0,
                    width, height, 1));

        glFinish();

        newData++;
        expectedData++;

        verify_and_update_kernel<<<(width*height)/256, 256>>>(writeSurface, readSurface, expectedData, newData, width, height);

        status = cuCtxSynchronize();
        if (status != CUDA_SUCCESS) {
            printf ("cuCtxSynchronize failed \n");
        }
    }

    cudaError_t err = cudaGetValueMismatch();
    if (err != cudaSuccess)
    {
        printf("Value mismatch seen when using CPU sync\n");
        cleanup(FAILURE);
    }

    // Clean up CUDA writeResource
    status = cuGraphicsUnregisterResource(writeResource);
    if (status != CUDA_SUCCESS) {
        printf("Failed to unregister %d", status);
        cleanup(FAILURE);
    }
    else {
        printf("Unregistered writeResource. \n");
    }

    // Clean up CUDA readResource
    status = cuGraphicsUnregisterResource(readResource);
    if (status != CUDA_SUCCESS) {
        printf("Failed to unregister %d", status);
        cleanup(FAILURE);
    }
    else {
        printf("Unregistered readResource. \n");
    }
}

/*
    Performs same function as checkSyncOnCPU
    Here instead of glFinish() and cuCtxSynchronize like in checkSyncOnCPU,
    we make use of EGLSync, CUDA Event and cuStreamWaitEvent, eglWaitSyncKHR to achieve the synchronization
    due to this CPU is not blocked for any synchronization needed between GL-EGL & CUDA operations
    all synchronizations happens on the GPU only.
*/
void checkSyncOnGPU(EGLDisplay dpy) 
{
    int z = 0;
    unsigned char expectedData, newData;
    cudaError_t err;
    CUresult status = CUDA_SUCCESS;
    CUstream stream;
    CUevent timingDisabledEvent;
    CUDA_RESOURCE_DESC wdsc,rdsc;
    memset(&wdsc, 0, sizeof(wdsc));
    memset(&rdsc, 0, sizeof(rdsc));

    expectedData = 0;
    newData = 1;

    wdsc.resType = CU_RESOURCE_TYPE_ARRAY;
    wdsc.res.array.hArray = writeArray;
    CUsurfObject writeSurface;
    rdsc.resType = CU_RESOURCE_TYPE_ARRAY;
    rdsc.res.array.hArray = readArray;
    CUsurfObject readSurface;

    status = cuSurfObjectCreate(&writeSurface, &wdsc);
    if (status != CUDA_SUCCESS)
    {
        printf("Surface bounding failed with status %d\n",status);
        cleanup(FAILURE);
    }
    status = cuSurfObjectCreate(&readSurface, &rdsc);
    if (status != CUDA_SUCCESS)
    {
        printf("Surface bounding failed\n");
        cleanup(FAILURE);
    }

    status = cuStreamCreate(&stream, CU_STREAM_DEFAULT);
    if(status != CUDA_SUCCESS) {
        printf("Stream creation failed\n");
        cleanup(FAILURE);
    }

    // Creates timing disabled event which uses non-blocking synchronization
    status = cuEventCreate(&timingDisabledEvent, CU_EVENT_DISABLE_TIMING);
    if(status != CUDA_SUCCESS) {
        printf("Default event creation failed\n");
        cleanup(FAILURE);
    }

    /*
        1. We perform texture-to-texture copy in GLES which is async function
        2. Followed by creating EGLSync and a CUDA Event from that EGLSync object
        3. Using cuStreamWaitEvent() we wait in GPU for the GLES to finish texture copy.
        4. CUDA kernel verfiy_and_update_kernel verifies if the copied data by GLES is correct,
           and it updates the buffer with new values.
        5. This is followed by eglWaitSyncKHR() which waits for the cuda kernel to finish,
           so that in the next iteration GLES can perform the copying of the updated buffer to write texture,
    */
    for(z = 0; z < itr; z++) {

        //GL call to copy from read texture to write texture
        GL_SAFE_CALL(glCopyImageSubData(tex[GL_READ], GL_TEXTURE_2D, 0, 0, 0, 0,
                    tex[GL_WRITE], GL_TEXTURE_2D, 0, 0, 0, 0,
                    width, height, 1));

        EGLSyncKHR eglSyncForGL, eglSyncForCuda;
        EGLBoolean egl_status = EGL_TRUE;
        EGLAttribKHR eglattrib[] = { EGL_CUDA_EVENT_HANDLE_NV, (EGLAttrib)timingDisabledEvent, EGL_NONE};

        CUevent cudaEGLSyncEvent;

        eglSyncForGL = eglCreateSyncKHR(dpy, EGL_SYNC_FENCE_KHR, NULL);

        if(eglSyncForGL == EGL_NO_SYNC_KHR) {
            printf(" EGL Sync creation failed\n");
            cleanup(FAILURE);
        }

        status = cuEventCreateFromEGLSync(&cudaEGLSyncEvent, eglSyncForGL, CU_EVENT_DEFAULT);
        if(status != CUDA_SUCCESS) {
            printf("CUDA event creation from EGLSync failed\n");
            cleanup(FAILURE);
        }

        // We wait from CUDA in GPU for GL-EGL operation completion
        status = cuStreamWaitEvent(stream, cudaEGLSyncEvent, 0);
        if(status != CUDA_SUCCESS) {
            printf("Stream wait for event created from EGLSync failed\n");
            cleanup(FAILURE);
        }

        egl_status = eglDestroySyncKHR(dpy, eglSyncForGL);
        if(egl_status != EGL_TRUE) {
            printf("EGL sync object destruction failed\n");
            cleanup(FAILURE);
        }

        newData++;
        expectedData++;

        // Verifies the values in readSurface which is copied by glCopyImageSubData()
        // And writes value of newData into writeSurface
        verify_and_update_kernel<<<(width*height)/256, 256, 0, stream>>>(writeSurface,readSurface, expectedData, newData, width, height);

        status  = cuEventDestroy(cudaEGLSyncEvent);
        if(status != CUDA_SUCCESS) {
            printf("Event Destroy failed\n");
            cleanup(FAILURE);
        }

        status = cuEventRecord(timingDisabledEvent, stream);
        if(status != CUDA_SUCCESS) {
            printf("Event Record failed\n");
            cleanup(FAILURE);
        }

        // creating an EGL sync object linked to a CUDA event object
        eglSyncForCuda = eglCreateSync64KHR(dpy, EGL_SYNC_CUDA_EVENT_NV, eglattrib);

        // We wait from EGL for CUDA operation completion
        egl_status = eglWaitSyncKHR(dpy, eglSyncForCuda, 0);
        if(egl_status != EGL_TRUE) {
            printf("eglWaitSyncKHR failed\n");
            cleanup(FAILURE);
        }
        egl_status = eglDestroySyncKHR(dpy, eglSyncForCuda);
        if(egl_status != EGL_TRUE) {
            printf("EGL sync object destruction failed\n");
            cleanup(FAILURE);
        }
    }

    err = cudaGetValueMismatch();
    if (err != cudaSuccess)
    {
        printf("Value mismatch seen when using GPU sync\n");
        cleanup(FAILURE);
    }

    // Clean up CUDA writeResource
    status = cuGraphicsUnregisterResource(writeResource);
    if (status != CUDA_SUCCESS) {
        printf("Failed to unregister %d", status);
        cleanup(FAILURE);
    }
    else {
        printf("Unregistered writeResource. \n");
    }

    // Clean up CUDA readResource
    status = cuGraphicsUnregisterResource(readResource);
    if (status != CUDA_SUCCESS) {
        printf("Failed to unregister %d", status);
        cleanup(FAILURE);
    }
    else {
        printf("Unregistered readResource. \n");
    }
}

// Verifies the values in readSurface whether they are expected ones
// And writes value of newData into writeSurface
__global__ void verify_and_update_kernel(CUsurfObject write, CUsurfObject read, char expected, char newval, int width, int height)
{
    unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;

    if ( x < width && y < height) 
    {
        uchar4 check;
        surf2Dread(&check, read, x * 4, y);
        if (check.x != expected || check.y != expected || check.z != expected || check.w != expected)
        {
            printf("Mismatch found in values read[0]= %u read[1]= %u read[2]= %u read[3]= %u expected is %u\n",check.x, check.y, check.z, check.w, expected);
            numErrors++;
            return;
        }
        uchar4 data = make_uchar4(newval, newval, newval, newval);
        surf2Dwrite(data, write, x * 4, y);
    }
}

__global__ void getNumErrors(int *numErr) {
    *numErr = numErrors;
}

extern "C" cudaError_t cudaGetValueMismatch()
{
    int numErr_h;
    int *numErr_d = NULL;
    cudaError_t err = cudaSuccess;

    err = cudaMalloc(&numErr_d, sizeof(int));
    if (err != cudaSuccess) {
        printf("Cuda Main: cudaMemcpy failed with %s\n",cudaGetErrorString(err));
        cudaFree(numErr_d);
        return err;
    }

    getNumErrors<<<1,1>>>(numErr_d);
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("Cuda Main: cudaDeviceSynchronize failed with %s\n",cudaGetErrorString(err));
    }
    err = cudaMemcpy(&numErr_h, numErr_d, sizeof(int), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        printf("Cuda Main: cudaMemcpy failed with %s\n",cudaGetErrorString(err));
        cudaFree(numErr_d);
        return err;
    }
    err = cudaFree(numErr_d);
    if (err != cudaSuccess) {
        printf("Cuda Main: cudaFree failed with %s\n",cudaGetErrorString(err));
        return err;
    }
    if (numErr_h > 0) {
        return cudaErrorUnknown;
    }
    return cudaSuccess;
}


// Clean up state and exit. If status is SUCCESS, regression success is printed
// to stdout. This will happen if the glut timer is triggered. If status is
// anything else, the regression failure message is printed.
void cleanup(int status)
{
    GLenum glErr = GL_NO_ERROR;
    cudaError cudaErr = cudaSuccess;
    int exitStatus = status;

    // Clean up GL
    if(*tex)
    {
        GL_SAFE_CALL_NO_CLEANUP(glDeleteTextures(2, tex), glErr);
    }

    // Print test status and exit
    if(glErr != GL_NO_ERROR || cudaErr != cudaSuccess)
        exitStatus = FAILURE;

    printStatus(exitStatus);

    cleanExit = 1;

    graphics_close_window();

    if(exitStatus == FAILURE)
        exit(EXIT_FAILURE);

    exit(0);
}


void exitHandler(void)
{
    if (!cleanExit) {
        printf("&&&& EGLSync_CUDAEvent_Interop unexpected failure \n");
        printStatus(FAILURE);
    }
}

// Print test success or fail for regression testing
void printStatus(int status)
{
    switch (status) {
        case SUCCESS:
            printf("&&&& EGLSync_CUDAEvent_Interop PASSED\n");
            break;
        case WAIVED:
            printf("&&&& EGLSync_CUDAEvent_Interop WAIVED\n");
            break;
        default:
            printf("&&&& EGLSync_CUDAEvent_Interop FAILED\n");
            break;
    }
    fflush(stdout);
}