 	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,32,32,32,cudaChannelFormatKindFloat);
 	cudaExtent extent = make_cudaExtent(N , N, N);
 	checkCudaErrors(cudaMalloc3DArray(&dataArray, &channelDesc,extent));
+
+	dataTex.filterMode = cudaFilterModeLinear;
 	//Generate data
 	float4*** hostvfield; 
 //	hostvfield = (float4*) malloc(N*N*N*sizeof(float4));

	std::cout << 5 << std::endl;
	checkCudaErrors(cudaBindTextureToArray(dataTex, dataArray, channelDesc));
	float4 *d_lines, *h_lines;
//	double time = 3.141592653*2.0;
	int steps = N;
	int cores = 1;
	int blocks = 1;
	float dt = spacing;
	float4 startloc = {origin,0,0,0};
	std::cout << hostvfield[127][127][127].x << std::endl;
	checkCudaErrors(cudaMalloc(&d_lines,blocks*cores*steps*sizeof(float4)));

	h_lines = (float4*) malloc(blocks*cores*steps*sizeof(float4));
	RK4line<<<cores,blocks>>>(d_lines, dt, steps, startloc);
	checkCudaErrors(cudaMemcpy(h_lines, d_lines, blocks*cores*steps*sizeof(float4), cudaMemcpyDeviceToHost));
	for(unsigned int i=0; i<steps; i++) {
		std::cout << "x= " << h_lines[i].x << "; y= "<< h_lines[i].y << " "<< h_lines[i].x*h_lines[i].x+h_lines[i].y*h_lines[i].y;
	}
	free(hostvfield[0][0]);
	free(hostvfield[0]);
