Commit a2c755c88e2930d26bbb3796639e408c1cac1228
1 parent
db331d8a
start new network class
Showing
4 changed files
with
53 additions
and
11 deletions
Show diff stats
1 | +# -*- coding: utf-8 -*- | ||
2 | +""" | ||
3 | +Created on Sat Jan 19 2018 | ||
4 | + | ||
5 | +@author: Jiabing | ||
6 | +""" | ||
7 | + | ||
8 | +import struct | ||
9 | +import numpy as np | ||
10 | +import scipy as sp | ||
11 | +import networkx as nx | ||
12 | +import matplotlib.pyplot as plt | ||
13 | +import math | ||
14 | + | ||
15 | +class Point: | ||
16 | + def __init__(self, x, y, z, radius): | ||
17 | + self.x = x | ||
18 | + self.y = y | ||
19 | + self.z = z | ||
20 | + self.r = radius | ||
21 | + | ||
22 | + | ||
23 | +class Fiber: | ||
24 | + def __init__(self, fiber_idx,point_idx): | ||
25 | + self.fidx = fiber_idx | ||
26 | + self.pidx = point_idx | ||
27 | + | ||
28 | + | ||
29 | + | ||
30 | +class Node: | ||
31 | + def __init__(self, point_idx, fidx): | ||
32 | + self.pidx= point_idx | ||
33 | + self.fidx = fidx | ||
34 | + | ||
35 | + | ||
36 | +#class NWT: | ||
37 | + | ||
38 | +#class network(point, Fiber, Node): | ||
0 | \ No newline at end of file | 39 | \ No newline at end of file |
stim/iVote/ivote2.cuh
@@ -152,13 +152,17 @@ namespace stim { | @@ -152,13 +152,17 @@ namespace stim { | ||
152 | 152 | ||
153 | 153 | ||
154 | template<typename T> | 154 | template<typename T> |
155 | - void cpu_ivote2(T* cpuI, unsigned int rmax, size_t x, size_t y, bool invert = false, T t = 0, int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) { | 155 | + void cpu_ivote2(T* cpuI, unsigned int rmax, size_t x, size_t y, float &gpu_time, bool invert = false, T t = 0, int iter = 8, T phi = 15.0f * (float)stim::PI / 180, int conn = 8, bool debug = false) { |
156 | size_t bytes = x*y * sizeof(T); | 156 | size_t bytes = x*y * sizeof(T); |
157 | T* gpuI; //allocate space on the gpu to save the input image | 157 | T* gpuI; //allocate space on the gpu to save the input image |
158 | + | ||
159 | + gpuTimer_start(); | ||
158 | HANDLE_ERROR(cudaMalloc(&gpuI, bytes)); | 160 | HANDLE_ERROR(cudaMalloc(&gpuI, bytes)); |
159 | HANDLE_ERROR(cudaMemcpy(gpuI, cpuI, bytes, cudaMemcpyHostToDevice)); //copy the image to the gpu | 161 | HANDLE_ERROR(cudaMemcpy(gpuI, cpuI, bytes, cudaMemcpyHostToDevice)); //copy the image to the gpu |
160 | stim::gpu_ivote2<T>(gpuI, rmax, x, y, invert, t, iter, phi, conn, debug); //call the gpu version of the ivote | 162 | stim::gpu_ivote2<T>(gpuI, rmax, x, y, invert, t, iter, phi, conn, debug); //call the gpu version of the ivote |
161 | HANDLE_ERROR(cudaMemcpy(cpuI, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output to the cpu | 163 | HANDLE_ERROR(cudaMemcpy(cpuI, gpuI, bytes, cudaMemcpyDeviceToHost)); //copy the output to the cpu |
164 | + | ||
165 | + gpu_time = gpuTimer_end(); | ||
162 | } | 166 | } |
163 | } | 167 | } |
164 | #endif | 168 | #endif |
165 | \ No newline at end of file | 169 | \ No newline at end of file |
stim/iVote/ivote2/local_max.cuh
@@ -10,7 +10,7 @@ namespace stim{ | @@ -10,7 +10,7 @@ namespace stim{ | ||
10 | 10 | ||
11 | // this kernel calculates the local maximum for finding the cell centers | 11 | // this kernel calculates the local maximum for finding the cell centers |
12 | template<typename T> | 12 | template<typename T> |
13 | - __global__ void cuda_local_max(T* gpuCenters, T* gpuVote, int conn, int x, int y){ | 13 | + __global__ void cuda_local_max(T* gpuCenters, T* gpuVote, int conn, int x, int y){ |
14 | 14 | ||
15 | // calculate the 2D coordinates for this current thread. | 15 | // calculate the 2D coordinates for this current thread. |
16 | int xi = blockIdx.x * blockDim.x + threadIdx.x; | 16 | int xi = blockIdx.x * blockDim.x + threadIdx.x; |
@@ -20,16 +20,16 @@ namespace stim{ | @@ -20,16 +20,16 @@ namespace stim{ | ||
20 | return; | 20 | return; |
21 | 21 | ||
22 | // convert 2D coordinates to 1D | 22 | // convert 2D coordinates to 1D |
23 | - int i = yi * x + xi; | 23 | + int i = yi * x + xi; |
24 | 24 | ||
25 | gpuCenters[i] = 0; //initialize the value at this location to zero | 25 | gpuCenters[i] = 0; //initialize the value at this location to zero |
26 | 26 | ||
27 | T val = gpuVote[i]; | 27 | T val = gpuVote[i]; |
28 | 28 | ||
29 | - for(int xl = xi - conn; xl < xi + conn; xl++){ | ||
30 | - for(int yl = yi - conn; yl < yi + conn; yl++){ | 29 | + for(unsigned int xl = xi - conn; xl < xi + conn; xl++){ |
30 | + for(unsigned int yl = yi - conn; yl < yi + conn; yl++){ | ||
31 | if(xl >= 0 && xl < x && yl >= 0 && yl < y){ | 31 | if(xl >= 0 && xl < x && yl >= 0 && yl < y){ |
32 | - int il = yl * x + xl; | 32 | + unsigned int il = yl * x + xl; |
33 | if(gpuVote[il] > val){ | 33 | if(gpuVote[il] > val){ |
34 | return; | 34 | return; |
35 | } | 35 | } |
@@ -47,7 +47,7 @@ namespace stim{ | @@ -47,7 +47,7 @@ namespace stim{ | ||
47 | } | 47 | } |
48 | 48 | ||
49 | template<typename T> | 49 | template<typename T> |
50 | - void gpu_local_max(T* gpuCenters, T* gpuVote, unsigned int conn, size_t x, size_t y){ | 50 | + void gpu_local_max(T* gpuCenters, T* gpuVote, unsigned int conn, unsigned int x, unsigned int y){ |
51 | 51 | ||
52 | unsigned int max_threads = stim::maxThreadsPerBlock(); | 52 | unsigned int max_threads = stim::maxThreadsPerBlock(); |
53 | /*dim3 threads(max_threads, 1); | 53 | /*dim3 threads(max_threads, 1); |
stim/image/image.h
@@ -275,10 +275,10 @@ public: | @@ -275,10 +275,10 @@ public: | ||
275 | int channels = cvImage.channels(); | 275 | int channels = cvImage.channels(); |
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; | ||
279 | - //if (C() == 1) //if this is a single-color image, just copy the data | ||
280 | - // type_copy<T, T>(cv_ptr, img, size()); | ||
281 | - memcpy(img, cv_ptr, bytes()); | 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()); | ||
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()); |
284 | #else | 284 | #else |