#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define WIDTH 2
#define HEIGHT 2

texture<float, 2> texRef;
float* d_tex;

__global__ void kernel(float* d_tex)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    float val = tex2D(texRef,x,y);
    if(val < 0)
        d_tex[offset] = 1;

}

int main(void)
{

    int imageSize = WIDTH*HEIGHT*sizeof(float);
    cudaMalloc((void**)&d_tex, imageSize);
    const cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D(NULL, texRef, d_tex, desc, WIDTH, HEIGHT, sizeof(float)*WIDTH);

    float h_tex[4];

    float* tempTex = (float*)malloc(imageSize);

    /*
    [ 1 1
      1 1 ]
    */

    tempTex[0] = -1; tempTex[1] = -1; tempTex[2] = -1;
    tempTex[3] = -1;

    /* tempTex
    [ 1 1
      1 1 ]
    */

    cudaMemcpy(d_tex, tempTex,imageSize,cudaMemcpyHostToDevice);
    free(tempTex);

    dim3    blocks(1,1);
    dim3    threads(2,2);
    kernel<<<blocks,threads>>>(d_tex);

    cudaMemcpy( &h_tex, d_tex, imageSize, cudaMemcpyDeviceToHost );

    cudaUnbindTexture(texRef);
    cudaFree(d_tex);
    free(h_tex);

    return 0;
}