Commit ba75208e4845f04d46f273f9cde8bd42f5a2683b
1 parent
41acaf5d
change the name of mathvec.h to vector.h, fixed a bug in colormap.h
Showing
4 changed files
with
5 additions
and
146 deletions
Show diff stats
stim/cuda/imageproc/gaussian_blur.cuh deleted
1 | -#ifndef STIM_CUDA_GAUSSIAN_BLUR_H | ||
2 | -#define STIM_CUDA_GAUSSIAN_BLUR_H | ||
3 | - | ||
4 | -#include <iostream> | ||
5 | -#include <cuda.h> | ||
6 | -#include "gaussian_blur.cuh" | ||
7 | -#include <stim/cuda/devices.h> | ||
8 | - | ||
9 | -#define pi 3.14159 | ||
10 | - | ||
11 | -// This CUDA kernel applies a Gaussian blur along the x axis | ||
12 | -template<typename T> | ||
13 | -__global__ void gaussian_blur_x(T* out, T* in, T sigma, unsigned int x, unsigned int y){ | ||
14 | - | ||
15 | - //calculate the 1D image index for this thread | ||
16 | - int i = blockIdx.x * blockDim.x + threadIdx.x; | ||
17 | - | ||
18 | - //convert this to 2D pixel coordinates | ||
19 | - int yi = i / x; | ||
20 | - int xi = i - (yi * x); | ||
21 | - | ||
22 | - int k_size = sigma * 4; //calculate the kernel size | ||
23 | - | ||
24 | - int gx, gy, gi; //variables to store global coordinates | ||
25 | - T sum = 0; //variable stores the integral of the kernel times the image | ||
26 | - T G; //stores the value of the gaussian kernel | ||
27 | - | ||
28 | - out[i] = 0; | ||
29 | - //for each element of the kernel | ||
30 | - for(int ki = -k_size; ki <= k_size; ki++){ | ||
31 | - | ||
32 | - //calculate the gaussian value | ||
33 | - G = 1.0 / (sigma * sqrt(2 * pi)) * exp(-(ki*ki) / (2*sigma*sigma)); | ||
34 | - | ||
35 | - //calculate the global coordinates for this point in the kernel | ||
36 | - gx = xi + ki; | ||
37 | - gy = yi; | ||
38 | - | ||
39 | - //make sure we are inside the bounds of the image | ||
40 | - if(gx >= 0 && gy >= 0 && gx < x && gy < y){ | ||
41 | - gi = gy * x + gx; | ||
42 | - | ||
43 | - //perform the integration | ||
44 | - sum += G * in[gi]; | ||
45 | - } | ||
46 | - } | ||
47 | - | ||
48 | - //set the output value to the integral of the function times the kernel | ||
49 | - out[i] = sum; | ||
50 | -} | ||
51 | - | ||
52 | -// This CUDA kernel applies a Gaussian blur along the y axis. | ||
53 | -template<typename T> | ||
54 | -__global__ void gaussian_blur_y(T* out, T* in, T sigma, unsigned int x, unsigned int y){ | ||
55 | - | ||
56 | - //calculate the 1D image index for this thread | ||
57 | - int i = blockIdx.x * blockDim.x + threadIdx.x; | ||
58 | - | ||
59 | - //convert this to 2D pixel coordinates | ||
60 | - int yi = i / x; | ||
61 | - int xi = i - (yi * x); | ||
62 | - | ||
63 | - int k_size = sigma * 4; //calculate the kernel size | ||
64 | - | ||
65 | - int gx, gy, gi; //variables to store global coordinates | ||
66 | - T sum = 0; //variable stores the integral of the kernel times the image | ||
67 | - T G; //stores the value of the gaussian kernel | ||
68 | - | ||
69 | - out[i] = 0; | ||
70 | - //for each element of the kernel | ||
71 | - for(int ki = -k_size; ki <= k_size; ki++){ | ||
72 | - | ||
73 | - //calculate the gaussian value | ||
74 | - G = 1.0 / (sigma * sqrt(2 * pi)) * exp(-(ki*ki) / (2*sigma*sigma)); | ||
75 | - | ||
76 | - //calculate the global coordinates for this point in the kernel | ||
77 | - gx = xi; | ||
78 | - gy = yi + ki; | ||
79 | - | ||
80 | - //make sure we are inside the bounds of the image | ||
81 | - if(gx >= 0 && gy >= 0 && gx < x && gy < y){ | ||
82 | - gi = gy * x + gx; | ||
83 | - | ||
84 | - //perform the integration | ||
85 | - sum += G * in[gi]; | ||
86 | - } | ||
87 | - } | ||
88 | - | ||
89 | - //set the output value to the integral of the function times the kernel | ||
90 | - out[i] = sum; | ||
91 | -} | ||
92 | - | ||
93 | -/// Applies a Gaussian blur to a 2D image stored on the GPU | ||
94 | -template<typename T> | ||
95 | -void gpu_gaussian_blur_2d(T* image, T sigma, unsigned int x, unsigned int y){ | ||
96 | - | ||
97 | - //get the number of pixels in the image | ||
98 | - unsigned int pixels = x * y; | ||
99 | - unsigned int bytes = sizeof(T) * pixels; | ||
100 | - | ||
101 | - //allocate temporary space on the GPU | ||
102 | - T* gpuI1; | ||
103 | - cudaMalloc(&gpuI1, bytes); | ||
104 | - | ||
105 | - //get the maximum number of threads per block for the CUDA device | ||
106 | - int threads = stim::maxThreadsPerBlock(); | ||
107 | - | ||
108 | - //calculate the number of blocks | ||
109 | - int blocks = pixels / threads + (pixels%threads == 0 ? 0:1); | ||
110 | - | ||
111 | - //run the kernel for the x dimension | ||
112 | - gaussian_blur_x <<< blocks, threads >>>(gpuI1, image, sigma, x, y); | ||
113 | - | ||
114 | - //run the kernel for the y dimension | ||
115 | - gaussian_blur_y <<< blocks, threads >>>(image, gpuI1, sigma, x, y); | ||
116 | - | ||
117 | -} | ||
118 | - | ||
119 | -/// Applies a Gaussian blur to a 2D image stored on the CPU | ||
120 | -template<typename T> | ||
121 | -void cpu_gaussian_blur_2d(T* image, T sigma, unsigned int x, unsigned int y){ | ||
122 | - | ||
123 | - //get the number of pixels in the image | ||
124 | - unsigned int pixels = x * y; | ||
125 | - unsigned int bytes = sizeof(T) * pixels; | ||
126 | - | ||
127 | - //allocate space on the GPU | ||
128 | - T* gpuI0; | ||
129 | - cudaMalloc(&gpuI0, bytes); | ||
130 | - | ||
131 | - //copy the image data to the GPU | ||
132 | - cudaMemcpy(gpuI0, image, bytes, cudaMemcpyHostToDevice); | ||
133 | - | ||
134 | - //run the GPU-based version of the algorithm | ||
135 | - gpu_gaussian_blur_2d<T>(gpuI0, sigma, x, y); | ||
136 | - | ||
137 | - //copy the image data from the device | ||
138 | - cudaMemcpy(image, gpuI0, bytes, cudaMemcpyDeviceToHost); | ||
139 | -} | ||
140 | - | ||
141 | -#endif |
stim/envi/binary.h
@@ -100,7 +100,7 @@ public: | @@ -100,7 +100,7 @@ public: | ||
100 | /// @param filename is the name of the binary file | 100 | /// @param filename is the name of the binary file |
101 | /// @param r is a STIM vector specifying the size of the binary file along each dimension | 101 | /// @param r is a STIM vector specifying the size of the binary file along each dimension |
102 | /// @param h is the length (in bytes) of any header file (default zero) | 102 | /// @param h is the length (in bytes) of any header file (default zero) |
103 | - bool open(std::string filename, vec<unsigned int, D> r, unsigned int h = 0){ | 103 | + bool open(std::string filename, vec<unsigned int> r, unsigned int h = 0){ |
104 | 104 | ||
105 | for(unsigned int i = 0; i < D; i++) //set the dimensions of the binary file object | 105 | for(unsigned int i = 0; i < D; i++) //set the dimensions of the binary file object |
106 | R[i] = r[i]; | 106 | R[i] = r[i]; |
@@ -117,7 +117,7 @@ public: | @@ -117,7 +117,7 @@ public: | ||
117 | /// @param filename is the name of the binary file to be created | 117 | /// @param filename is the name of the binary file to be created |
118 | /// @param r is a STIM vector specifying the size of the file along each dimension | 118 | /// @param r is a STIM vector specifying the size of the file along each dimension |
119 | /// @offset specifies how many bytes to offset the file (used to leave room for a header) | 119 | /// @offset specifies how many bytes to offset the file (used to leave room for a header) |
120 | - bool create(std::string filename, vec<unsigned int, D> r, unsigned int offset = 0){ | 120 | + bool create(std::string filename, vec<unsigned int> r, unsigned int offset = 0){ |
121 | 121 | ||
122 | std::ofstream target(filename.c_str(), std::ios::binary); | 122 | std::ofstream target(filename.c_str(), std::ios::binary); |
123 | 123 |
stim/math/mathvec.h renamed to stim/math/vector.h
stim/visualization/colormap.h
@@ -287,7 +287,7 @@ static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned int nVals, T | @@ -287,7 +287,7 @@ static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned int nVals, T | ||
287 | } | 287 | } |
288 | 288 | ||
289 | template<class T> | 289 | template<class T> |
290 | -static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned int nVals, colormapType cm = cmGrayscale)//, bool positive = false) | 290 | +static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned int nVals, colormapType cm = cmGrayscale) |
291 | { | 291 | { |
292 | //computes the max and min range automatically | 292 | //computes the max and min range automatically |
293 | 293 | ||
@@ -329,13 +329,13 @@ static void cpu2image(T* cpuSource, std::string fileDest, unsigned int x_size, u | @@ -329,13 +329,13 @@ static void cpu2image(T* cpuSource, std::string fileDest, unsigned int x_size, u | ||
329 | } | 329 | } |
330 | 330 | ||
331 | template<typename T> | 331 | template<typename T> |
332 | -static void cpu2image(T* cpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, colormapType cm = cmGrayscale, bool positive = false) | 332 | +static void cpu2image(T* cpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, colormapType cm = cmGrayscale) |
333 | { | 333 | { |
334 | //allocate a color buffer | 334 | //allocate a color buffer |
335 | unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size); | 335 | unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size); |
336 | 336 | ||
337 | //do the mapping | 337 | //do the mapping |
338 | - cpu2cpu<T>(cpuSource, cpuBuffer, x_size * y_size, cm, positive); | 338 | + cpu2cpu<T>(cpuSource, cpuBuffer, x_size * y_size, cm); |
339 | 339 | ||
340 | //copy the buffer to an image | 340 | //copy the buffer to an image |
341 | buffer2image(cpuBuffer, fileDest, x_size, y_size); | 341 | buffer2image(cpuBuffer, fileDest, x_size, y_size); |