8a86bd56
David Mayerich
changed rts names...
|
1
2
|
#ifndef STIM_COLORMAP_H
#define STIM_COLORMAP_H
|
7006df5f
David Mayerich
reformat of direc...
|
3
4
5
|
#include <string>
#include <stdlib.h>
|
a2bf1d08
David Mayerich
general bug fixes...
|
6
|
#include <cmath>
|
a2bf1d08
David Mayerich
general bug fixes...
|
7
8
9
10
11
|
#ifdef _WIN32
#include <float.h>
#endif
|
3eb12494
David Mayerich
removed CUDA depe...
|
12
|
#ifdef __CUDACC__
|
58e55e6b
David Mayerich
fixed colormap.h ...
|
13
|
#include "cublas_v2.h"
|
7d1d153a
Pavel Govyadinov
fixed the include...
|
14
|
#include <stim/cuda/cudatools/error.h>
|
3eb12494
David Mayerich
removed CUDA depe...
|
15
|
#endif
|
7006df5f
David Mayerich
reformat of direc...
|
16
|
|
76396b52
David Mayerich
removed Qt from t...
|
17
18
|
//saving an image to a file uses the CImg library
//this currently throws a lot of "unreachable" warnings (as of GCC 4.8.2, nvcc 6.5.12)
|
7d1d153a
Pavel Govyadinov
fixed the include...
|
19
|
#include <stim/image/image.h>
|
7006df5f
David Mayerich
reformat of direc...
|
20
|
|
7006df5f
David Mayerich
reformat of direc...
|
21
|
|
76396b52
David Mayerich
removed Qt from t...
|
22
|
#define BREWER_CTRL_PTS 11
|
7006df5f
David Mayerich
reformat of direc...
|
23
24
25
26
27
28
29
30
31
32
33
34
35
|
static float BREWERCP[BREWER_CTRL_PTS*4] = {0.192157f, 0.211765f, 0.584314f, 1.0f,
0.270588f, 0.458824f, 0.705882f, 1.0f,
0.454902f, 0.678431f, 0.819608f, 1.0f,
0.670588f, 0.85098f, 0.913725f, 1.0f,
0.878431f, 0.952941f, 0.972549f, 1.0f,
1.0f, 1.0f, 0.74902f, 1.0f,
0.996078f, 0.878431f, 0.564706f, 1.0f,
0.992157f, 0.682353f, 0.380392f, 1.0f,
0.956863f, 0.427451f, 0.262745f, 1.0f,
0.843137f, 0.188235f, 0.152941f, 1.0f,
0.647059f, 0.0f, 0.14902f, 1.0f};
|
4ea8daa8
Jiaming Guo
add necessary fun...
|
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
|
//static float BREWERCP[BREWER_CTRL_PTS * 4] = { 0.192157f, 0.584314f, 0.211765f, 1.0f,
// 0.270588f, 0.705882f, 0.458824f, 1.0f,
// 0.454902f, 0.819608f, 0.678431f, 1.0f,
// 0.670588f, 0.913725f, 0.85098f, 1.0f,
// 0.878431f, 0.972549f, 0.952941f, 1.0f,
// 1.0f, 0.74902f, 1.0f, 1.0f,
// 0.996078f, 0.878431f, 0.564706f, 1.0f,
// 0.992157f, 0.682353f, 0.380392f, 1.0f,
// 0.956863f, 0.427451f, 0.262745f, 1.0f,
// 0.843137f, 0.188235f, 0.152941f, 1.0f,
// 0.647059f, 0.0f, 0.14902f, 1.0f };
//static float BREWERCP[BREWER_CTRL_PTS * 4] = { 0.0f, 0.407843f, 0.215686f, 1.0f,
// 0.101960f, 0.596078f, 0.313725f, 1.0f,
// 0.4f, 0.741176f, 0.388235f, 1.0f,
// 0.650980f, 0.850980f, 0.415686f, 1.0f,
// 0.850980f, 0.937254f, 0.545098f, 1.0f,
// 1.0f, 1.0f, 0.749019f, 1.0f,
// 0.996078f, 0.878431f, 0.545098f, 1.0f,
// 0.992156f, 0.682352f, 0.380392f, 1.0f,
// 0.956862f, 0.427450f, 0.262745f, 1.0f,
// 0.843137f, 0.188235f, 0.152941f, 1.0f,
// 0.647058f, 0.0f, 0.149019f, 1.0f };
|
7006df5f
David Mayerich
reformat of direc...
|
60
61
62
63
64
65
|
#ifdef __CUDACC__
texture<float4, cudaTextureType1D> cudaTexBrewer;
static cudaArray* gpuBrewer;
#endif
|
8a86bd56
David Mayerich
changed rts names...
|
66
|
namespace stim{
|
7006df5f
David Mayerich
reformat of direc...
|
67
|
|
a9275be5
David Mayerich
added vector fiel...
|
68
|
enum colormapType {cmBrewer, cmGrayscale, cmRainbow};
|
7006df5f
David Mayerich
reformat of direc...
|
69
|
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
70
|
static void buffer2image(unsigned char* buffer, std::string filename, size_t width, size_t height)
|
7006df5f
David Mayerich
reformat of direc...
|
71
|
{
|
fbf0ab02
David Mayerich
added support for...
|
72
|
/*unsigned char* non_interleaved = (unsigned char*)malloc(x_size * y_size * 3);
|
76396b52
David Mayerich
removed Qt from t...
|
73
74
75
76
77
78
|
unsigned int S = x_size * y_size;
for(unsigned int i = 0; i < S; i++){
non_interleaved[i + 0 * S] = buffer[i * 3 + 0];
non_interleaved[i + 1 * S] = buffer[i * 3 + 1];
non_interleaved[i + 2 * S] = buffer[i * 3 + 2];
|
fbf0ab02
David Mayerich
added support for...
|
79
|
}*/
|
76396b52
David Mayerich
removed Qt from t...
|
80
81
|
//create an image object
|
fbf0ab02
David Mayerich
added support for...
|
82
83
84
|
//cimg_library::CImg<unsigned char> image(non_interleaved, x_size, y_size, 1, 3);
//image.save(filename.c_str());
image<unsigned char> I;
|
c8c976a9
David Mayerich
replaced CImg wit...
|
85
|
I.set_interleaved_rgb(buffer, width, height);
|
fbf0ab02
David Mayerich
added support for...
|
86
|
I.save(filename);
|
7006df5f
David Mayerich
reformat of direc...
|
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
|
}
#ifdef __CUDACC__
static void initBrewer()
{
//initialize the Brewer colormap
//allocate CPU space
float4 cpuColorMap[BREWER_CTRL_PTS];
//define control rtsPoints
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));
HANDLE_ERROR(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));
}
static void destroyBrewer()
{
HANDLE_ERROR(cudaFreeArray(gpuBrewer));
|
7006df5f
David Mayerich
reformat of direc...
|
136
137
138
139
140
141
142
143
144
145
146
147
|
}
template<class T>
__global__ static void applyBrewer(T* gpuSource, unsigned char* gpuDest, unsigned int N, T minVal = 0, T maxVal = 1)
{
int i = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
if(i >= N) return;
//compute the normalized value on [minVal maxVal]
float a = (gpuSource[i] - minVal) / (maxVal - minVal);
|
8d4f0940
David Mayerich
ERROR in plane wa...
|
148
149
150
|
//compensate for the additional space at the edges
a *= (T)(BREWER_CTRL_PTS - 1)/(T)(BREWER_CTRL_PTS);
|
7006df5f
David Mayerich
reformat of direc...
|
151
|
//lookup the color
|
8d4f0940
David Mayerich
ERROR in plane wa...
|
152
|
float shift = (T)1/(2*BREWER_CTRL_PTS);
|
7006df5f
David Mayerich
reformat of direc...
|
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
|
float4 color = tex1D(cudaTexBrewer, a+shift);
//float4 color = tex1D(cudaTexBrewer, a);
gpuDest[i * 3 + 0] = 255 * color.x;
gpuDest[i * 3 + 1] = 255 * color.y;
gpuDest[i * 3 + 2] = 255 * color.z;
}
template<class T>
__global__ static void applyGrayscale(T* gpuSource, unsigned char* gpuDest, unsigned int N, T minVal = 0, T maxVal = 1)
{
int i = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
if(i >= N) return;
//compute the normalized value on [minVal maxVal]
float a = (gpuSource[i] - minVal) / (maxVal - minVal);
//threshold
|
a9275be5
David Mayerich
added vector fiel...
|
171
172
173
174
|
if(a > 1)
a = 1;
if(a < 0)
a = 0;
|
7006df5f
David Mayerich
reformat of direc...
|
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
|
gpuDest[i * 3 + 0] = 255 * a;
gpuDest[i * 3 + 1] = 255 * a;
gpuDest[i * 3 + 2] = 255 * a;
}
template<class T>
static void gpu2gpu(T* gpuSource, unsigned char* gpuDest, unsigned int nVals, T minVal = 0, T maxVal = 1, colormapType cm = cmGrayscale, int blockDim = 128)
{
//This function converts a scalar field on the GPU to a color image on the GPU
int gridX = (nVals + blockDim - 1)/blockDim;
int gridY = 1;
if(gridX > 65535)
{
gridY = (gridX + 65535 - 1) / 65535;
gridX = 65535;
}
dim3 dimGrid(gridX, gridY);
|
7006df5f
David Mayerich
reformat of direc...
|
193
194
195
196
197
198
|
if(cm == cmGrayscale)
applyGrayscale<<<dimGrid, blockDim>>>(gpuSource, gpuDest, nVals, minVal, maxVal);
else if(cm == cmBrewer)
{
initBrewer();
applyBrewer<<<dimGrid, blockDim>>>(gpuSource, gpuDest, nVals, minVal, maxVal);
|
7006df5f
David Mayerich
reformat of direc...
|
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
|
destroyBrewer();
}
}
template<class T>
static void gpu2cpu(T* gpuSource, unsigned char* cpuDest, unsigned int nVals, T minVal, T maxVal, colormapType cm = cmGrayscale)
{
//this function converts a scalar field on the GPU to a color image on the CPU
//first create the color image on the GPU
//allocate GPU memory for the color image
unsigned char* gpuDest;
HANDLE_ERROR(cudaMalloc( (void**)&gpuDest, sizeof(unsigned char) * nVals * 3 ));
|
7006df5f
David Mayerich
reformat of direc...
|
215
216
|
//create the image on the gpu
gpu2gpu(gpuSource, gpuDest, nVals, minVal, maxVal, cm);
|
654c8934
David Mayerich
CUDA_CALLABLE fix...
|
217
|
|
7006df5f
David Mayerich
reformat of direc...
|
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
|
//copy the image from the GPU to the CPU
HANDLE_ERROR(cudaMemcpy(cpuDest, gpuDest, sizeof(unsigned char) * nVals * 3, cudaMemcpyDeviceToHost));
HANDLE_ERROR(cudaFree( gpuDest ));
}
template<typename T>
static void gpu2image(T* gpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, T valMin, T valMax, colormapType cm = cmGrayscale)
{
//allocate a color buffer
unsigned char* cpuBuffer = NULL;
cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size);
//do the mapping
gpu2cpu<T>(gpuSource, cpuBuffer, x_size * y_size, valMin, valMax, cm);
//copy the buffer to an image
buffer2image(cpuBuffer, fileDest, x_size, y_size);
free(cpuBuffer);
}
|
4252d827
David Mayerich
ivote3 fixes and ...
|
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
|
/// save a GPU image to a file using automatic scaling
template<typename T>
static void gpu2image(T* gpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, colormapType cm = cmGrayscale){
size_t N = x_size * y_size; //calculate the total number of elements in the image
cublasStatus_t stat;
cublasHandle_t handle;
stat = cublasCreate(&handle); //create a cuBLAS handle
if (stat != CUBLAS_STATUS_SUCCESS){ //test for failure
printf ("CUBLAS initialization failed\n");
exit(1);
}
int i_min, i_max;
stat = cublasIsamin(handle, (int)N, gpuSource, 1, &i_min);
if (stat != CUBLAS_STATUS_SUCCESS){ //test for failure
printf ("CUBLAS Error: failed to calculate minimum r value.\n");
exit(1);
}
stat = cublasIsamax(handle, (int)N, gpuSource, 1, &i_max);
if (stat != CUBLAS_STATUS_SUCCESS){ //test for failure
printf ("CUBLAS Error: failed to calculate maximum r value.\n");
exit(1);
}
cublasDestroy(handle);
i_min--; //cuBLAS uses 1-based indexing for Fortran compatibility
i_max--;
T v_min, v_max; //allocate space to store the minimum and maximum values
HANDLE_ERROR( cudaMemcpy(&v_min, gpuSource + i_min, sizeof(T), cudaMemcpyDeviceToHost) ); //copy the min and max values from the device to the CPU
HANDLE_ERROR( cudaMemcpy(&v_max, gpuSource + i_max, sizeof(T), cudaMemcpyDeviceToHost) );
|
3669ed62
David Mayerich
fixed fused multi...
|
274
275
276
|
gpu2image<T>(gpuSource, fileDest, x_size, y_size, min(v_min, v_max), max(v_min, v_max), cm);
|
4252d827
David Mayerich
ivote3 fixes and ...
|
277
278
|
}
|
7006df5f
David Mayerich
reformat of direc...
|
279
280
281
|
#endif
template<class T>
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
282
|
static void cpuApplyBrewer(T* cpuSource, unsigned char* cpuDest, size_t N, T minVal = 0, T maxVal = 1)
|
7006df5f
David Mayerich
reformat of direc...
|
283
|
{
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
284
|
for(size_t i=0; i<N; i++)
|
7006df5f
David Mayerich
reformat of direc...
|
285
286
|
{
//compute the normalized value on [minVal maxVal]
|
27b826a8
David Mayerich
added a spherical...
|
287
288
289
290
291
|
float a;
if(minVal != maxVal)
a = (cpuSource[i] - minVal) / (maxVal - minVal);
else
a = 0.5;
|
a2bf1d08
David Mayerich
general bug fixes...
|
292
293
294
295
296
297
298
|
#ifdef _WIN32
if(!_finite(a)) a = 1; //deal with infinite and NaN values (return maximum in all cases)
#else
if(!std::isfinite(a)) a = 1;
#endif
else if(a < 0) a = 0;
else if(a > 1) a = 1;
|
7006df5f
David Mayerich
reformat of direc...
|
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
|
float c = a * (float)(BREWER_CTRL_PTS-1);
int ptLow = (int)c;
float m = c - (float)ptLow;
//std::cout<<m<<std::endl;
float r, g, b;
if(ptLow == BREWER_CTRL_PTS - 1)
{
r = BREWERCP[ptLow * 4 + 0];
g = BREWERCP[ptLow * 4 + 1];
b = BREWERCP[ptLow * 4 + 2];
}
else
{
|
a9275be5
David Mayerich
added vector fiel...
|
314
315
316
|
r = BREWERCP[ptLow * 4 + 0] * (1-m) + BREWERCP[ (ptLow+1) * 4 + 0] * m;
g = BREWERCP[ptLow * 4 + 1] * (1-m) + BREWERCP[ (ptLow+1) * 4 + 1] * m;
b = BREWERCP[ptLow * 4 + 2] * (1-m) + BREWERCP[ (ptLow+1) * 4 + 2] * m;
|
7006df5f
David Mayerich
reformat of direc...
|
317
318
319
|
}
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
320
321
322
|
cpuDest[i * 3 + 0] = (unsigned char)(255 * r);
cpuDest[i * 3 + 1] = (unsigned char)(255 * g);
cpuDest[i * 3 + 2] = (unsigned char)(255 * b);
|
7006df5f
David Mayerich
reformat of direc...
|
323
324
325
326
327
|
}
}
template<class T>
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
328
|
static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, size_t nVals, T valMin, T valMax, colormapType cm = cmGrayscale)
|
7006df5f
David Mayerich
reformat of direc...
|
329
330
331
332
333
334
335
336
337
|
{
if(cm == cmBrewer)
cpuApplyBrewer(cpuSource, cpuDest, nVals, valMin, valMax);
else if(cm == cmGrayscale)
{
int i;
float a;
float range = valMax - valMin;
|
27b826a8
David Mayerich
added a spherical...
|
338
|
|
7006df5f
David Mayerich
reformat of direc...
|
339
340
341
|
for(i = 0; i<nVals; i++)
{
//normalize to the range [valMin valMax]
|
27b826a8
David Mayerich
added a spherical...
|
342
343
344
345
346
|
if(range != 0)
a = (cpuSource[i] - valMin) / range;
else
a = 0.5;
|
a9275be5
David Mayerich
added vector fiel...
|
347
348
|
if(a < 0) a = 0;
if(a > 1) a = 1;
|
7006df5f
David Mayerich
reformat of direc...
|
349
|
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
350
351
352
|
cpuDest[i * 3 + 0] = (unsigned char)(255 * a);
cpuDest[i * 3 + 1] = (unsigned char)(255 * a);
cpuDest[i * 3 + 2] = (unsigned char)(255 * a);
|
7006df5f
David Mayerich
reformat of direc...
|
353
354
355
356
357
|
}
}
}
template<class T>
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
358
|
static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned long long nVals, colormapType cm = cmGrayscale)
|
7006df5f
David Mayerich
reformat of direc...
|
359
360
361
362
|
{
//computes the max and min range automatically
//find the largest magnitude value
|
27b826a8
David Mayerich
added a spherical...
|
363
364
365
|
T maxVal = cpuSource[0];
T minVal = cpuSource[0];
for(int i=1; i<nVals; i++)
|
7006df5f
David Mayerich
reformat of direc...
|
366
|
{
|
27b826a8
David Mayerich
added a spherical...
|
367
368
369
370
|
if(cpuSource[i] > maxVal)
maxVal = cpuSource[i];
if(cpuSource[i] < minVal)
minVal = cpuSource[i];
|
7006df5f
David Mayerich
reformat of direc...
|
371
372
|
}
|
27b826a8
David Mayerich
added a spherical...
|
373
|
cpu2cpu(cpuSource, cpuDest, nVals, minVal, maxVal, cm);
|
7006df5f
David Mayerich
reformat of direc...
|
374
375
376
377
378
379
|
}
template<typename T>
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
380
|
static void cpu2image(T* cpuSource, std::string fileDest, size_t x_size, size_t y_size, T valMin, T valMax, colormapType cm = cmGrayscale)
|
7006df5f
David Mayerich
reformat of direc...
|
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
|
{
//allocate a color buffer
unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size);
//do the mapping
cpu2cpu<T>(cpuSource, cpuBuffer, x_size * y_size, valMin, valMax, cm);
//copy the buffer to an image
buffer2image(cpuBuffer, fileDest, x_size, y_size);
free(cpuBuffer);
}
template<typename T>
|
9d3ba0b1
David Mayerich
added stim::hsi a...
|
396
|
static void cpu2image(T* cpuSource, std::string fileDest, size_t x_size, size_t y_size, colormapType cm = cmGrayscale)
|
7006df5f
David Mayerich
reformat of direc...
|
397
398
399
400
401
|
{
//allocate a color buffer
unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size);
//do the mapping
|
ba75208e
David Mayerich
change the name o...
|
402
|
cpu2cpu<T>(cpuSource, cpuBuffer, x_size * y_size, cm);
|
7006df5f
David Mayerich
reformat of direc...
|
403
404
405
406
407
408
409
410
411
412
413
|
//copy the buffer to an image
buffer2image(cpuBuffer, fileDest, x_size, y_size);
free(cpuBuffer);
}
} //end namespace colormap and rts
#endif
|