f1402849
 dmayerich
renewed commit
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
|
#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));
}
|
