#include "rtsVector3d.h"
#include "cudaHandleError.h"

#define BREWER_CTRL_PTS 11
texture<float4, cudaTextureType1D> cudaTexBrewer;
cudaArray* gpuBrewer;


__device__ float4 colorBrewer(float val, float min_range, float max_range, float min_fade = 0.000)
{
	float t = (val - min_range)/(max_range - min_range)*((float)(BREWER_CTRL_PTS-2)/BREWER_CTRL_PTS);
	float shift = 1.0/BREWER_CTRL_PTS;
	float4 color = tex1D(cudaTexBrewer, t+shift);
	if(t < min_fade)
		color.w = (val - min_range)/(max_range - min_range)/min_fade;
	return color;
}


void rts_cudaInitBrewer()
{
	//allocate CPU space
	float4 cpuColorMap[BREWER_CTRL_PTS];

	//define control points
	cpuColorMap[0] = make_float4(0.192157f, 0.211765f, 0.584314f, 1.0f);
	cpuColorMap[1] = make_float4(0.270588f, 0.458824f, 0.705882f, 1.0f);
	cpuColorMap[2] = make_float4(0.454902f, 0.678431f, 0.819608f, 1.0f);
	cpuColorMap[3] = make_float4(0.670588f, 0.85098f, 0.913725f, 1.0f);
	cpuColorMap[4] = make_float4(0.878431f, 0.952941f, 0.972549f, 1.0f);
	cpuColorMap[5] = make_float4(1.0f, 1.0f, 0.74902f, 1.0f);
	cpuColorMap[6] = make_float4(0.996078f, 0.878431f, 0.564706f, 1.0f);
	cpuColorMap[7] = make_float4(0.992157f, 0.682353f, 0.380392f, 1.0f);
	cpuColorMap[8] = make_float4(0.956863f, 0.427451f, 0.262745f, 1.0f);
	cpuColorMap[9] = make_float4(0.843137f, 0.188235f, 0.152941f, 1.0f);
	cpuColorMap[10] = make_float4(0.647059f, 0.0f, 0.14902f, 1.0f);


	int width = BREWER_CTRL_PTS;
	int height = 0;


	// allocate array and copy colormap data
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);

	HANDLE_ERROR(cudaMallocArray(&gpuBrewer, &channelDesc, width, height));

	cudaMemcpyToArray(gpuBrewer, 0, 0, cpuColorMap, sizeof(float4)*width, cudaMemcpyHostToDevice);

	// set texture parameters
    cudaTexBrewer.addressMode[0] = cudaAddressModeClamp;
	//texBrewer.addressMode[1] = cudaAddressModeClamp;
    cudaTexBrewer.filterMode = cudaFilterModeLinear;
    cudaTexBrewer.normalized = true;  // access with normalized texture coordinates

	// Bind the array to the texture
    HANDLE_ERROR( cudaBindTextureToArray( cudaTexBrewer, gpuBrewer, channelDesc));

}