#include <helper_cuda.h>
#include <helper_functions.h>
#include <curand.h>

//static texture<float, cudaTextureType3D, cudaReadModeElementType> NoiseTextures[4];//texture Array
static texture<float, cudaTextureType3D, cudaReadModeElementType> NoiseTexture;
float *d_NoiseTest;//Device Array with random floats
int SizeNoiseTest = 32;
int sqrSizeNoiseTest = 32768;

void CreateTexture();

__global__ void AccesTexture()
{
        float test = tex3D(NoiseTexture,threadIdx.x,threadIdx.y,threadIdx.z);//by using this the error occurs
        printf("thread: %d,%d,%d, value: %f\n", threadIdx.x, threadIdx.y, threadIdx.z, test);
}

int
main(int argc, char **argv)
{
        CreateTexture();
}
void CreateTexture()
{
        //curand Random Generator (needs compiler link -lcurand)
        curandGenerator_t gen;
        cudaMalloc((void **)&d_NoiseTest, sqrSizeNoiseTest*sizeof(float));//Allocation of device Array
        curandCreateGenerator(&gen,CURAND_RNG_PSEUDO_DEFAULT);
        curandSetPseudoRandomGeneratorSeed(gen,1235ULL);
        curandGenerateUniform(gen, d_NoiseTest, sqrSizeNoiseTest);//writing data to d_NoiseTest
        curandDestroyGenerator(gen);

        //cudaArray Descriptor
        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
        //cuda Array
        cudaArray *d_cuArr;
        checkCudaErrors(cudaMalloc3DArray(&d_cuArr, &channelDesc, make_cudaExtent(SizeNoiseTest*sizeof(float),SizeNoiseTest,SizeNoiseTest), 0));
        cudaMemcpy3DParms copyParams = {0};


                //Array creation
                copyParams.srcPtr   = make_cudaPitchedPtr(d_NoiseTest, SizeNoiseTest*sizeof(float), SizeNoiseTest, SizeNoiseTest);
                copyParams.dstArray = d_cuArr;
                copyParams.extent   = make_cudaExtent(SizeNoiseTest,SizeNoiseTest,SizeNoiseTest);
                copyParams.kind     = cudaMemcpyDeviceToDevice;
                checkCudaErrors(cudaMemcpy3D(&copyParams));
                //Array creation End

                //new Bind
                // set texture parameters
                NoiseTexture.normalized = false;                      // use non-normalized texture coordinates
                NoiseTexture.filterMode = cudaFilterModeLinear;      // linear interpolation
                NoiseTexture.addressMode[0] = cudaAddressModeWrap;   // wrap texture coordinates
                NoiseTexture.addressMode[1] = cudaAddressModeWrap;
                NoiseTexture.addressMode[2] = cudaAddressModeWrap;

                // bind array to 3D texture
                checkCudaErrors(cudaBindTextureToArray(NoiseTexture, d_cuArr, channelDesc));
                //end Bind
//      cudaFreeArray(d_cuArr);
        checkCudaErrors(cudaPeekAtLastError());
        AccesTexture<<<1,dim3(4,4,4)>>>();
        checkCudaErrors(cudaPeekAtLastError());
        cudaDeviceSynchronize();
        checkCudaErrors(cudaPeekAtLastError());
}