Commit 8da0df3ee72d360aef8d26bc1186853477c76cbd
1 parent
0954c632
whatever
Showing
4 changed files
with
23 additions
and
21 deletions
Show diff stats
python/network.py
@@ -318,14 +318,14 @@ class Network: | @@ -318,14 +318,14 @@ class Network: | ||
318 | R = (200, 200, 200) | 318 | R = (200, 200, 200) |
319 | 319 | ||
320 | #generate a meshgrid of the appropriate size and resolution to surround the network | 320 | #generate a meshgrid of the appropriate size and resolution to surround the network |
321 | - lower, upper = self.aabb(self.N, self.F) #get the space occupied by the network | 321 | + lower, upper = self.aabb() #get the space occupied by the network |
322 | x = np.linspace(lower[0], upper[0], R[0]) #get the grid points for uniform sampling of this space | 322 | x = np.linspace(lower[0], upper[0], R[0]) #get the grid points for uniform sampling of this space |
323 | y = np.linspace(lower[1], upper[1], R[1]) | 323 | y = np.linspace(lower[1], upper[1], R[1]) |
324 | z = np.linspace(lower[2], upper[2], R[2]) | 324 | z = np.linspace(lower[2], upper[2], R[2]) |
325 | X, Y, Z = np.meshgrid(x, y, z) | 325 | X, Y, Z = np.meshgrid(x, y, z) |
326 | #Z = 150 * numpy.ones(X.shape) | 326 | #Z = 150 * numpy.ones(X.shape) |
327 | 327 | ||
328 | - | 328 | + |
329 | Q = np.stack((X, Y, Z), 3) | 329 | Q = np.stack((X, Y, Z), 3) |
330 | 330 | ||
331 | 331 |
python/network_dpy.py
stim/cuda/templates/gaussian_blur.cuh
@@ -9,34 +9,34 @@ | @@ -9,34 +9,34 @@ | ||
9 | 9 | ||
10 | 10 | ||
11 | namespace stim{ | 11 | namespace stim{ |
12 | - namespace cuda{ | 12 | + namespace cuda { |
13 | 13 | ||
14 | template<typename T> | 14 | template<typename T> |
15 | - void gen_gaussian(T* out, T sigma, unsigned int width){ | 15 | + void gen_gaussian(T* out, T sigma, unsigned int width) { |
16 | 16 | ||
17 | //fill the kernel with a gaussian | 17 | //fill the kernel with a gaussian |
18 | - for(unsigned int xi = 0; xi < width; xi++){ | 18 | + for (unsigned int xi = 0; xi < width; xi++) { |
19 | 19 | ||
20 | - float x = (float)xi - (float)(width/2); //calculate the x position of the gaussian | ||
21 | - float g = 1.0 / (sigma * sqrt(2 * 3.14159)) * exp( - (x*x) / (2*sigma*sigma) ); | 20 | + float x = (float)xi - (float)(width / 2); //calculate the x position of the gaussian |
21 | + float g = 1.0 / (sigma * sqrt(2 * 3.14159)) * exp(-(x*x) / (2 * sigma*sigma)); | ||
22 | out[xi] = g; | 22 | out[xi] = g; |
23 | } | 23 | } |
24 | 24 | ||
25 | } | 25 | } |
26 | 26 | ||
27 | template<typename T> | 27 | template<typename T> |
28 | - void tex_gaussian_blur2(T* out, T sigma, unsigned int x, unsigned int y, cudaTextureObject_t texObj, cudaArray* cuArray){ | 28 | + void tex_gaussian_blur2(T* out, T sigma, unsigned int x, unsigned int y, cudaTextureObject_t texObj, cudaArray* cuArray) { |
29 | 29 | ||
30 | //allocate space for the kernel | 30 | //allocate space for the kernel |
31 | unsigned int kwidth = sigma * 8 + 1; | 31 | unsigned int kwidth = sigma * 8 + 1; |
32 | - float* kernel0 = (float*) malloc( kwidth * sizeof(float) ); | 32 | + float* kernel0 = (float*)malloc(kwidth * sizeof(float)); |
33 | 33 | ||
34 | //fill the kernel with a gaussian | 34 | //fill the kernel with a gaussian |
35 | gen_gaussian(kernel0, sigma, kwidth); | 35 | gen_gaussian(kernel0, sigma, kwidth); |
36 | 36 | ||
37 | //copy the kernel to the GPU | 37 | //copy the kernel to the GPU |
38 | T* gpuKernel0; | 38 | T* gpuKernel0; |
39 | - HANDLE_ERROR(cudaMalloc(&gpuKernel0, kwidth*sizeof(T))); | 39 | + HANDLE_ERROR(cudaMalloc(&gpuKernel0, kwidth * sizeof(T))); |
40 | HANDLE_ERROR(cudaMemcpy(gpuKernel0, kernel0, kwidth * sizeof(T), cudaMemcpyHostToDevice)); | 40 | HANDLE_ERROR(cudaMemcpy(gpuKernel0, kernel0, kwidth * sizeof(T), cudaMemcpyHostToDevice)); |
41 | 41 | ||
42 | //perform the gaussian blur as a separable convolution | 42 | //perform the gaussian blur as a separable convolution |
@@ -48,18 +48,18 @@ namespace stim{ | @@ -48,18 +48,18 @@ namespace stim{ | ||
48 | } | 48 | } |
49 | 49 | ||
50 | template<typename T> | 50 | template<typename T> |
51 | - void gpu_gaussian_blur2(T* image, T sigma, unsigned int x, unsigned int y){ | 51 | + void gpu_gaussian_blur2(T* image, T sigma, unsigned int x, unsigned int y) { |
52 | 52 | ||
53 | //allocate space for the kernel | 53 | //allocate space for the kernel |
54 | unsigned int kwidth = sigma * 8 + 1; | 54 | unsigned int kwidth = sigma * 8 + 1; |
55 | - float* kernel0 = (float*) malloc( kwidth * sizeof(float) ); | 55 | + float* kernel0 = (float*)malloc(kwidth * sizeof(float)); |
56 | 56 | ||
57 | //fill the kernel with a gaussian | 57 | //fill the kernel with a gaussian |
58 | gen_gaussian(kernel0, sigma, kwidth); | 58 | gen_gaussian(kernel0, sigma, kwidth); |
59 | 59 | ||
60 | //copy the kernel to the GPU | 60 | //copy the kernel to the GPU |
61 | T* gpuKernel0; | 61 | T* gpuKernel0; |
62 | - HANDLE_ERROR(cudaMalloc(&gpuKernel0, kwidth*sizeof(T))); | 62 | + HANDLE_ERROR(cudaMalloc(&gpuKernel0, kwidth * sizeof(T))); |
63 | HANDLE_ERROR(cudaMemcpy(gpuKernel0, kernel0, kwidth * sizeof(T), cudaMemcpyHostToDevice)); | 63 | HANDLE_ERROR(cudaMemcpy(gpuKernel0, kernel0, kwidth * sizeof(T), cudaMemcpyHostToDevice)); |
64 | 64 | ||
65 | //perform the gaussian blur as a separable convolution | 65 | //perform the gaussian blur as a separable convolution |
@@ -71,21 +71,24 @@ namespace stim{ | @@ -71,21 +71,24 @@ namespace stim{ | ||
71 | 71 | ||
72 | /// Applies a Gaussian blur to a 2D image stored on the CPU | 72 | /// Applies a Gaussian blur to a 2D image stored on the CPU |
73 | template<typename T> | 73 | template<typename T> |
74 | - void cpu_gaussian_blur2(T* image, T sigma, unsigned int x, unsigned int y){ | 74 | + void cpu_gaussian_blur2(T* image, T sigma, unsigned int x, unsigned int y, float &gpu_time) { |
75 | 75 | ||
76 | + gpuTimer_start(); | ||
76 | //allocate space for the kernel | 77 | //allocate space for the kernel |
77 | unsigned int kwidth = sigma * 8 + 1; | 78 | unsigned int kwidth = sigma * 8 + 1; |
78 | - float* kernel0 = (float*) malloc( kwidth * sizeof(float) ); | 79 | + float* kernel0 = (float*)malloc(kwidth * sizeof(float)); |
79 | 80 | ||
80 | //fill the kernel with a gaussian | 81 | //fill the kernel with a gaussian |
81 | gen_gaussian(kernel0, sigma, kwidth); | 82 | gen_gaussian(kernel0, sigma, kwidth); |
82 | 83 | ||
83 | //perform the gaussian blur as a separable convolution | 84 | //perform the gaussian blur as a separable convolution |
84 | stim::cuda::cpu_conv2sep<float>(image, x, y, kernel0, kwidth, kernel0, kwidth); | 85 | stim::cuda::cpu_conv2sep<float>(image, x, y, kernel0, kwidth, kernel0, kwidth); |
85 | - | 86 | + gpu_time = gpuTimer_end(); |
87 | + | ||
86 | } | 88 | } |
89 | + | ||
87 | 90 | ||
88 | - }; | ||
89 | -}; | 91 | + } |
92 | +} | ||
90 | 93 | ||
91 | #endif | 94 | #endif |
stim/image/image.h
@@ -276,8 +276,8 @@ public: | @@ -276,8 +276,8 @@ public: | ||
276 | allocate(cols, rows, channels); //allocate space for the image | 276 | allocate(cols, rows, channels); //allocate space for the image |
277 | size_t img_bytes = bytes(); | 277 | size_t img_bytes = bytes(); |
278 | unsigned char* cv_ptr = (unsigned char*)cvImage.data; | 278 | unsigned char* cv_ptr = (unsigned char*)cvImage.data; |
279 | - if (C() == 1) //if this is a single-color image, just copy the data | ||
280 | - type_copy<unsigned char, T>(cv_ptr, img, size()); | 279 | + //if (C() == 1) //if this is a single-color image, just copy the data |
280 | + type_copy<unsigned char, T>(cv_ptr, img, size()); | ||
281 | //memcpy(img, cv_ptr, bytes()); | 281 | //memcpy(img, cv_ptr, bytes()); |
282 | if(C() == 3) //if this is a 3-color image, OpenCV uses BGR interleaving | 282 | if(C() == 3) //if this is a 3-color image, OpenCV uses BGR interleaving |
283 | from_opencv(cv_ptr, X(), Y()); | 283 | from_opencv(cv_ptr, X(), Y()); |