Commit c0e09133e61049c0068a3b3bac016ffe46e04360

Authored by Pavel Govyadinov
1 parent f2c883d8

STABLE: made template size a variable, changed the gl_spider and cuda code to support that.

stim/cuda/cuda_texture.cuh
@@ -134,6 +134,9 @@ namespace stim @@ -134,6 +134,9 @@ namespace stim
134 HANDLE_ERROR( 134 HANDLE_ERROR(
135 cudaDestroyTextureObject(tObj) 135 cudaDestroyTextureObject(tObj)
136 ); 136 );
  137 +// HANDLE_ERROR(
  138 +// cudaFreeArray(srcArray)
  139 +// );
137 } 140 }
138 141
139 //-------------------------------------------------------------------------// 142 //-------------------------------------------------------------------------//
stim/cuda/spider_cost.cuh
@@ -17,7 +17,7 @@ namespace stim{ @@ -17,7 +17,7 @@ namespace stim{
17 17
18 stim::cuda::cuda_texture t; //texture object. 18 stim::cuda::cuda_texture t; //texture object.
19 float* result; 19 float* result;
20 - float* print; 20 +// float* print;
21 21
22 ///Initialization function, allocates the memory and passes the necessary 22 ///Initialization function, allocates the memory and passes the necessary
23 ///handles from OpenGL and Cuda. 23 ///handles from OpenGL and Cuda.
@@ -42,7 +42,9 @@ namespace stim{ @@ -42,7 +42,9 @@ namespace stim{
42 __device__ 42 __device__
43 float Template(int x) 43 float Template(int x)
44 { 44 {
45 - if(x < 16/6 || x > 16*5/6 || (x > 16*2/6 && x < 16*4/6)){ 45 + if(x < 32/6 || x > 32*5/6 || (x > 32*2/6 && x < 32*4/6)){
  46 +// if(x < 2 || x > 13 || (x > 5 && x < 10)){
  47 +// if(x < ceilf(16/6) || x > floorf(16*5/6) || (x > floorf(16*2/6) && x < ceilf(16*4/6))){
46 return 1.0; 48 return 1.0;
47 }else{ 49 }else{
48 return 0.0; 50 return 0.0;
@@ -57,14 +59,15 @@ namespace stim{ @@ -57,14 +59,15 @@ namespace stim{
57 ///@param float* result --a pointer to the memory that stores the result. 59 ///@param float* result --a pointer to the memory that stores the result.
58 __global__ 60 __global__
59 //void get_diff (float *result) 61 //void get_diff (float *result)
60 - void get_diff (cudaTextureObject_t texIn, float *result) 62 + void get_diff (cudaTextureObject_t texIn, float *result, int dx, int dy)
61 { 63 {
62 - __shared__ float shared[16][8]; 64 +// __shared__ float shared[32][16];
  65 + extern __shared__ float shared[];
63 int x = threadIdx.x + blockIdx.x * blockDim.x; 66 int x = threadIdx.x + blockIdx.x * blockDim.x;
64 int y = threadIdx.y + blockIdx.y * blockDim.y; 67 int y = threadIdx.y + blockIdx.y * blockDim.y;
65 int x_t = threadIdx.x; 68 int x_t = threadIdx.x;
66 int y_t = threadIdx.y; 69 int y_t = threadIdx.y;
67 -// int idx = y*16+x; 70 + int idx = y_t*dx+x_t;
68 int g_idx = blockIdx.y; 71 int g_idx = blockIdx.y;
69 72
70 float valIn = tex2D<unsigned char>(texIn, x, y)/255.0; 73 float valIn = tex2D<unsigned char>(texIn, x, y)/255.0;
@@ -72,7 +75,7 @@ namespace stim{ @@ -72,7 +75,7 @@ namespace stim{
72 75
73 // print[idx] = abs(valIn); ///temporary 76 // print[idx] = abs(valIn); ///temporary
74 77
75 - shared[x_t][y_t] = abs(valIn-valTemp); 78 + shared[idx] = abs(valIn-valTemp);
76 79
77 __syncthreads(); 80 __syncthreads();
78 81
@@ -81,7 +84,8 @@ namespace stim{ @@ -81,7 +84,8 @@ namespace stim{
81 __syncthreads(); 84 __syncthreads();
82 if (x_t < step) 85 if (x_t < step)
83 { 86 {
84 - shared[x_t][y_t] += shared[x_t + step][y_t]; 87 +// shared[x_t][y_t] += shared[x_t + step][y_t];
  88 + shared[idx] += shared[y_t*dx+x_t+step];
85 } 89 }
86 __syncthreads(); 90 __syncthreads();
87 } 91 }
@@ -92,13 +96,14 @@ namespace stim{ @@ -92,13 +96,14 @@ namespace stim{
92 __syncthreads(); 96 __syncthreads();
93 if(y_t < step) 97 if(y_t < step)
94 { 98 {
95 - shared[x_t][y_t] += shared[x_t][y_t + step]; 99 +// shared[x_t][y_t] += shared[x_t][y_t + step];
  100 + shared[idx] += shared[(y_t+step)*dx+x_t];
96 } 101 }
97 __syncthreads(); 102 __syncthreads();
98 } 103 }
99 __syncthreads(); 104 __syncthreads();
100 if(x_t == 0 && y_t == 0) 105 if(x_t == 0 && y_t == 0)
101 - result[g_idx] = shared[0][0]; 106 + result[g_idx] = shared[0];
102 107
103 108
104 // //result[idx] = abs(valIn); 109 // //result[idx] = abs(valIn);
@@ -112,7 +117,7 @@ namespace stim{ @@ -112,7 +117,7 @@ namespace stim{
112 /// may work with other gl texture types, but untested. 117 /// may work with other gl texture types, but untested.
113 ///@param DIM_Y, the number of samples in the template. 118 ///@param DIM_Y, the number of samples in the template.
114 extern "C" 119 extern "C"
115 - stim::vec<int> get_cost(GLint texbufferID, GLenum texType, int DIM_Y) 120 + stim::vec<int> get_cost(GLint texbufferID, GLenum texType, int DIM_Y,int dx = 16, int dy = 8)
116 { 121 {
117 122
118 //Bind the Texture in GL and allow access to cuda. 123 //Bind the Texture in GL and allow access to cuda.
@@ -132,10 +137,10 @@ namespace stim{ @@ -132,10 +137,10 @@ namespace stim{
132 137
133 //cuda launch variables. 138 //cuda launch variables.
134 dim3 numBlocks(1, DIM_Y); 139 dim3 numBlocks(1, DIM_Y);
135 - dim3 threadsPerBlock(16, 8); 140 + dim3 threadsPerBlock(dx, dy);
136 141
137 142
138 - get_diff <<< numBlocks, threadsPerBlock >>> (t.getTexture(), result); 143 + get_diff <<< numBlocks, threadsPerBlock, dx*dy*sizeof(float) >>> (t.getTexture(), result, dx, dy);
139 144
140 HANDLE_ERROR( 145 HANDLE_ERROR(
141 cudaMemcpy(output, result, DIM_Y*sizeof(float), cudaMemcpyDeviceToHost) 146 cudaMemcpy(output, result, DIM_Y*sizeof(float), cudaMemcpyDeviceToHost)
@@ -155,6 +160,7 @@ namespace stim{ @@ -155,6 +160,7 @@ namespace stim{
155 t.UnmapCudaTexture(); 160 t.UnmapCudaTexture();
156 cleanUP(); 161 cleanUP();
157 ret[0] = idx; ret[1] = (int) output[idx]; 162 ret[0] = idx; ret[1] = (int) output[idx];
  163 +// std::cout << "The cost is " << output[idx] << std::endl;
158 free(output); 164 free(output);
159 return ret; 165 return ret;
160 } 166 }
stim/cuda/testKernel.cuh
@@ -29,7 +29,7 @@ @@ -29,7 +29,7 @@
29 __device__ 29 __device__
30 float templ(int x) 30 float templ(int x)
31 { 31 {
32 - if(x < 16/6 || x > 16*5/6 || (x > 16*2/6 && x < 16*4/6)){ 32 + if(x < 32/6 || x > 32*5/6 || (x > 32*2/6 && x < 32*4/6)){
33 return 1.0; 33 return 1.0;
34 }else{ 34 }else{
35 return 0.0; 35 return 0.0;
@@ -48,9 +48,9 @@ @@ -48,9 +48,9 @@
48 { 48 {
49 int x = threadIdx.x + blockIdx.x * blockDim.x; 49 int x = threadIdx.x + blockIdx.x * blockDim.x;
50 int y = threadIdx.y + blockIdx.y * blockDim.y; 50 int y = threadIdx.y + blockIdx.y * blockDim.y;
51 - int idx = y*64+x;  
52 -// int idx = y*32+x;  
53 -// int idx = y*16+x; 51 +// int idx = y*64+x;
  52 + int idx = y*32+x;
  53 + // int idx = y*16+x;
54 54
55 float valIn = tex2D<unsigned char>(texIn, x, y); 55 float valIn = tex2D<unsigned char>(texIn, x, y);
56 float templa = templ(x); 56 float templa = templ(x);
@@ -96,3 +96,33 @@ @@ -96,3 +96,33 @@
96 cleanUP(); 96 cleanUP();
97 } 97 }
98 98
  99 + void test(GLint texbufferID, GLenum texType, int x, int y, std::string nam)
  100 + {
  101 +
  102 + //Bind the Texture in GL and allow access to cuda.
  103 + tx.MapCudaTexture(texbufferID, texType);
  104 +
  105 + //initialize the return arrays.
  106 +
  107 + initArray(x,y);
  108 + dim3 numBlocks(1, y);
  109 + dim3 threadsPerBlock(x, 1);
  110 + int max_threads = stim::maxThreadsPerBlock();
  111 + //dim3 threads(max_threads, 1);
  112 + //dim3 blocks(x / threads.x + 1, y);
  113 + //dim3 numBlocks(2, 2);
  114 + //dim3 threadsPerBlock(8, 108);
  115 +
  116 +
  117 +// get_diff <<< blocks, threads >>> (tx.getTexture(), print);
  118 + get_diff <<< numBlocks, threadsPerBlock >>> (tx.getTexture(), print);
  119 +
  120 + cudaDeviceSynchronize();
  121 + stringstream name; //for debugging
  122 + name << nam.c_str();
  123 + stim::gpu2image<float>(print, name.str(),x,y,0,255);
  124 +
  125 + tx.UnmapCudaTexture();
  126 + cleanUP();
  127 + }
  128 +
stim/gl/gl_spider.h
@@ -24,7 +24,7 @@ @@ -24,7 +24,7 @@
24 #include "../../../volume-spider/fiber.h" 24 #include "../../../volume-spider/fiber.h"
25 #include "../../../volume-spider/glnetwork.h" 25 #include "../../../volume-spider/glnetwork.h"
26 #include <stim/visualization/cylinder.h> 26 #include <stim/visualization/cylinder.h>
27 -//#include <stim/cuda/testKernel.cuh> 27 +#include <stim/cuda/testKernel.cuh>
28 28
29 //#include <stim/cuda/testKernel.cuh> 29 //#include <stim/cuda/testKernel.cuh>
30 30
@@ -103,6 +103,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -103,6 +103,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
103 stim::vec<float> ps; 103 stim::vec<float> ps;
104 stim::vec<float> ups; 104 stim::vec<float> ups;
105 stim::vec<float> ds; 105 stim::vec<float> ds;
  106 +
  107 + static const float t_length = 16.0;
106 108
107 109
108 //--------------------------------------------------------------------------// 110 //--------------------------------------------------------------------------//
@@ -117,9 +119,12 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -117,9 +119,12 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
117 { 119 {
118 setMatrix(); //create the transformation matrix. 120 setMatrix(); //create the transformation matrix.
119 glCallList(dList); //move the templates to p, d, m. 121 glCallList(dList); //move the templates to p, d, m.
  122 + #ifdef TESTING
  123 + test(texbufferID, GL_TEXTURE_2D,2*t_length,numSamples*t_length, "Final_Cost_Direction.bmp");
  124 + #endif
120 int best = getCost(texbufferID,numSamples); //find min cost. 125 int best = getCost(texbufferID,numSamples); //find min cost.
121 stim::vec<float> next( //find next vector. 126 stim::vec<float> next( //find next vector.
122 - dV[best][0]*S[0]*R[0], 127 + dV[best][0]*S[0]*R[0],
123 dV[best][1]*S[1]*R[1], 128 dV[best][1]*S[1]*R[1],
124 dV[best][2]*S[2]*R[2], 129 dV[best][2]*S[2]*R[2],
125 0); 130 0);
@@ -139,6 +144,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -139,6 +144,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
139 { 144 {
140 setMatrix(); //create the transformation matrix. 145 setMatrix(); //create the transformation matrix.
141 glCallList(dList+1); //move the templates to p, d, m. 146 glCallList(dList+1); //move the templates to p, d, m.
  147 + #ifdef TESTING
  148 + test(ptexbufferID, GL_TEXTURE_2D,2*t_length, numSamplesPos*t_length, "Final_Cost_Position.bmp");
  149 + #endif
142 int best = getCost(ptexbufferID, numSamplesPos); //find min cost. 150 int best = getCost(ptexbufferID, numSamplesPos); //find min cost.
143 // std::cerr << best << std::endl; 151 // std::cerr << best << std::endl;
144 stim::vec<float> next( //find next position. 152 stim::vec<float> next( //find next position.
@@ -162,6 +170,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -162,6 +170,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
162 { 170 {
163 setMatrix(); //create the transformation. 171 setMatrix(); //create the transformation.
164 glCallList(dList+2); //move the templates to p, d, m. 172 glCallList(dList+2); //move the templates to p, d, m.
  173 + #ifdef TESTING
  174 + test(mtexbufferID, GL_TEXTURE_2D, 2*t_length, numSamplesMag*t_length, "Final_Cost_Position.bmp");
  175 + #endif
165 int best = getCost(mtexbufferID, numSamplesMag); //get best cost. 176 int best = getCost(mtexbufferID, numSamplesMag); //get best cost.
166 setMagnitude(m[0]*mV[best][0]); //adjust the magnitude. 177 setMagnitude(m[0]*mV[best][0]); //adjust the magnitude.
167 } 178 }
@@ -263,11 +274,90 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -263,11 +274,90 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
263 } 274 }
264 } 275 }
265 276
  277 +
  278 + float uniformRandom()
  279 + {
  280 + return ( (float)(rand()))/( (float)(RAND_MAX));
  281 + }
  282 +
  283 + float normalRandom()
  284 + {
  285 + float u1 = uniformRandom();
  286 + float u2 = uniformRandom();
  287 + return cos(2.0*atan(1.0)*u2)*sqrt(-1.0*log(u1));
  288 + }
  289 +
  290 + stim::vec<float> uniformRandVector()
  291 + {
  292 + stim::vec<float> r(uniformRandom(), uniformRandom(), 1.0);
  293 + return r;
  294 + }
  295 +
  296 + stim::vec<float> normalRandVector()
  297 + {
  298 + stim::vec<float> r(normalRandom(), normalRandom(), 1.0);
  299 + return r;
  300 + }
  301 +
266 302
267 //--------------------------------------------------------------------------// 303 //--------------------------------------------------------------------------//
268 //---------------------TEMPLATE CREATION METHODS----------------------------// 304 //---------------------TEMPLATE CREATION METHODS----------------------------//
269 //--------------------------------------------------------------------------// 305 //--------------------------------------------------------------------------//
270 306
  307 + void
  308 + MonteCarloDirectionVectors(int nSamples, float solidAngle = 2*M_PI)
  309 + {
  310 +///equal area projection.
  311 +// std::vector<stim::vec<float> > vecsUni;
  312 +// for(int i = 0; i < nSamples; i++)
  313 +// {
  314 +// stim::vec<float> a = uniformRandVector();
  315 +// a[1] = a[1]*2*M_PI;
  316 +// a[0] = a[0]*(solidAngle);
  317 +// a = a.cyl2cart();
  318 +// stim::vec<float> b(
  319 +// sqrt(1.-(pow(a[0],2)+(pow(a[1],2)))/4)*a[0],
  320 +// sqrt(1.-(pow(a[0],2)+(pow(a[1],2)))/4)*a[1],
  321 +// 1.-(pow(a[0],2)+pow(a[1],2))/2);
  322 +// 2.*a[0]/(1.+pow(a[0],2)+pow(a[1],2)),
  323 +// 2.*a[1]/(1.+pow(a[0],2)+pow(a[1],2)),
  324 +// (-1.0+pow(a[0],2)+pow(a[1],2))/(1.0+pow(a[0],2)+pow(a[1],2)));
  325 +// b = b.norm();
  326 +// b[2] = -b[2]; //this doeswn't want to change
  327 +// vecsUni.push_back(b);
  328 +// }
  329 + float PHI[2], Z[2], range;
  330 + PHI[0] = asin(solidAngle/2);
  331 + PHI[1] = asin(0);
  332 +
  333 + Z[0] = cos(PHI[0]);
  334 + Z[1] = cos(PHI[1]);
  335 +
  336 +
  337 + range = Z[0] - Z[1];
  338 +
  339 +
  340 + float z, theta, phi;
  341 +
  342 + std::vector<stim::vec<float> > vecsUni;
  343 + for(int i = 0; i < numSamplesPos; i++)
  344 + {
  345 + stim::vec<float> a(uniformRandom()*0.8, uniformRandom()*0.8, 0.0);
  346 + a[0] = a[0]-0.4;
  347 + a[1] = a[1]-0.4;
  348 + vecsUni.push_back(a);
  349 + }
  350 +
  351 + stringstream name;
  352 + for(int i = 0; i < numSamplesPos; i++)
  353 + name << vecsUni[i].str() << std::endl;
  354 +
  355 + std::ofstream outFile;
  356 + outFile.open("New_Pos_Vectors.txt");
  357 + outFile << name.str().c_str();
  358 + }
  359 +
  360 +
271 ///@param solidAngle, the size of the arc to sample. 361 ///@param solidAngle, the size of the arc to sample.
272 ///Method for populating the vector arrays with sampled vectors. 362 ///Method for populating the vector arrays with sampled vectors.
273 ///Objects created are rectangles the with the created directions. 363 ///Objects created are rectangles the with the created directions.
@@ -277,46 +367,45 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -277,46 +367,45 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
277 void 367 void
278 genDirectionVectors(float solidAngle = 5/M_PI*4) 368 genDirectionVectors(float solidAngle = 5/M_PI*4)
279 { 369 {
  370 +
280 //Set up the vectors necessary for Rectangle creation. 371 //Set up the vectors necessary for Rectangle creation.
281 vec<float> Y(1.0,0.0,0.0); //orthogonal vec. 372 vec<float> Y(1.0,0.0,0.0); //orthogonal vec.
282 vec<float> pos(0.0,0.0,0.0); 373 vec<float> pos(0.0,0.0,0.0);
283 vec<float> mag(1.0, 1.0, 1.0); 374 vec<float> mag(1.0, 1.0, 1.0);
284 vec<float> dir(0.0, 0.0, 1.0); 375 vec<float> dir(0.0, 0.0, 1.0);
285 376
286 - //Set up the variable necessary for vector creation.  
287 - vec<float> d_s = d.cart2sph().norm();  
288 - vec<float> temp(0,0,0);  
289 - int dim = (sqrt(numSamples)-1)/2;  
290 - float p0 = -M_PI; //phi angle in spherical coordinates.  
291 - float dt = solidAngle/(2.0 * ((float)dim + 1.0)); //step size in Theta.  
292 - float dp = p0/(2.0*((float)dim + 1.0)); //step size in Phi. 377 + float PHI[2], Z[2], range;
  378 + PHI[0] = solidAngle/2;
  379 + PHI[1] = asin(0);
293 380
  381 + Z[0] = cos(PHI[0]);
  382 + Z[1] = cos(PHI[1]);
  383 +
  384 + range = Z[0] - Z[1];
  385 +
  386 + float z, theta, phi;
294 glNewList(dList, GL_COMPILE); 387 glNewList(dList, GL_COMPILE);
295 - //Loop over the above defined space creating distinct vectors.  
296 - int idx = 0;  
297 - for(int i = -dim; i <= dim; i++){  
298 - for(int j = -dim; j <= dim; j++){  
299 -  
300 - //Create linear index  
301 - idx = (j+dim)+(i+dim)*((dim*2)+1);  
302 - temp[0] = d_s[0]; //rotate vector  
303 - temp[1] = d_s[1]+dp*(float) i;  
304 - temp[2] = d_s[2]+dt*(float) j;  
305 -  
306 - temp = (temp.sph2cart()).norm(); //back to cart  
307 - dV.push_back(temp);  
308 - if(cos(Y.dot(temp))< 0.087){ Y[0] = 0.0; Y[1] = 1.0;}  
309 - else{Y[0] = 1.0; Y[1] = 0.0;}  
310 -  
311 - hor = stim::rect<float>(mag,  
312 - pos, temp,  
313 - ((Y.cross(temp)).cross(temp)).norm());  
314 - ver = stim::rect<float>(mag,  
315 - pos, temp,  
316 - hor.n());  
317 - UpdateBuffer(0.0, 0.0+idx*8.0);  
318 - CHECK_OPENGL_ERROR 388 + for(int i = 0; i < numSamples; i++)
  389 + {
  390 + z = uniformRandom()*range + Z[1];
  391 + theta = uniformRandom()*2*M_PI;
  392 + phi = acos(z);
  393 + stim::vec<float> sph(1, theta, phi);
  394 + stim::vec<float> cart = sph.sph2cart();
  395 + dV.push_back(cart);
  396 + if(cos(Y.dot(cart)) < 0.087)
  397 + {
  398 + Y[0] = 0.0; Y[1] = 1.0;
  399 + }else{
  400 + Y[0] = 1.0; Y[1] = 0.0;
319 } 401 }
  402 + hor = stim::rect<float>(mag,
  403 + pos, cart,
  404 + ((Y.cross(cart)).cross(cart)).norm());
  405 + ver = stim::rect<float>(mag,
  406 + pos, cart,
  407 + hor.n());
  408 + UpdateBuffer(0.0, 0.0+i*t_length);
320 } 409 }
321 glEndList(); 410 glEndList();
322 } 411 }
@@ -337,36 +426,21 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -337,36 +426,21 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
337 vec<float> dir(0.0, 0.0, 1.0); 426 vec<float> dir(0.0, 0.0, 1.0);
338 427
339 //Set up the variable necessary for vector creation. 428 //Set up the variable necessary for vector creation.
340 - vec<float> temp(0,0,0);  
341 - int dim = (sqrt(numSamplesPos)-1)/2; //number of position vectors.  
342 - stim::rect<float> samplingPlane = //plane from which we pull position samples  
343 - stim::rect<float>(p, d);  
344 - samplingPlane.scale(mag[0]*delta, mag[0]*delta);  
345 - float step = 1.0/(dim); //step size.  
346 -  
347 - //Loop over the samples, keeping the original p samples in the center of the resulting texture to create a large number of position vectors.  
348 - int idx;  
349 glNewList(dList+1, GL_COMPILE); 429 glNewList(dList+1, GL_COMPILE);
350 - for(int i = -dim; i <= dim; i++){  
351 - for(int j = -dim; j <= dim; j++){  
352 - //Create linear index  
353 - idx = (j+dim)+(i+dim)*((dim*2)+1);  
354 -  
355 - temp = samplingPlane.p(  
356 - 0.5+step*i,  
357 - 0.5+step*j  
358 - );  
359 - pV.push_back(temp);  
360 - hor = stim::rect<float>(mag,  
361 - temp, dir,  
362 - ((Y.cross(d)).cross(d))  
363 - .norm());  
364 - ver = stim::rect<float>(mag,  
365 - temp, dir,  
366 - hor.n());  
367 - UpdateBuffer(0.0, 0.0+idx*8.0);  
368 - CHECK_OPENGL_ERROR  
369 - } 430 + for(int i = 0; i < numSamplesPos; i++)
  431 + {
  432 + stim::vec<float> temp = uniformRandVector();
  433 + temp = temp*delta*2.0 - delta/2.0;
  434 + temp[2] = 0.0;
  435 + pV.push_back(temp);
  436 + hor = stim::rect<float>(mag,
  437 + temp, dir,
  438 + ((Y.cross(d)).cross(d))
  439 + .norm());
  440 + ver = stim::rect<float>(mag,
  441 + temp, dir,
  442 + hor.n());
  443 + UpdateBuffer(0.0, 0.0+i*t_length);
370 } 444 }
371 glEndList(); 445 glEndList();
372 } 446 }
@@ -382,8 +456,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -382,8 +456,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
382 { 456 {
383 457
384 //Set up the vectors necessary for Rectangle creation. 458 //Set up the vectors necessary for Rectangle creation.
385 - vec<float> Y(1.0,0.0,0.0); //orthogonal vec.  
386 - vec<float> pos(0.0,0.0,0.0); 459 + vec<float> Y(1.0, 0.0, 0.0); //orthogonal vec.
  460 + vec<float> pos(0.0, 0.0, 0.0);
387 vec<float> mag(1.0, 1.0, 1.0); 461 vec<float> mag(1.0, 1.0, 1.0);
388 vec<float> dir(0.0, 0.0, 1.0); 462 vec<float> dir(0.0, 0.0, 1.0);
389 463
@@ -408,7 +482,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -408,7 +482,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
408 ver = stim::rect<float>(temp, 482 ver = stim::rect<float>(temp,
409 pos, dir, 483 pos, dir,
410 hor.n()); 484 hor.n());
411 - UpdateBuffer(0.0, 0.0+i*8.0); 485 + UpdateBuffer(0.0, 0.0+i*t_length);
412 CHECK_OPENGL_ERROR 486 CHECK_OPENGL_ERROR
413 } 487 }
414 glEndList(); 488 glEndList();
@@ -421,7 +495,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -421,7 +495,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
421 void 495 void
422 UpdateBuffer(float v_x, float v_y) 496 UpdateBuffer(float v_x, float v_y)
423 { 497 {
424 - float len = 8.0;  
425 stim::vec<float>p1; 498 stim::vec<float>p1;
426 stim::vec<float>p2; 499 stim::vec<float>p2;
427 stim::vec<float>p3; 500 stim::vec<float>p3;
@@ -442,19 +515,19 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -442,19 +515,19 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
442 p2[1], 515 p2[1],
443 p2[2] 516 p2[2]
444 ); 517 );
445 - glVertex2f(v_x+len, v_y); 518 + glVertex2f(v_x+t_length, v_y);
446 glTexCoord3f( 519 glTexCoord3f(
447 p3[0], 520 p3[0],
448 p3[1], 521 p3[1],
449 p3[2] 522 p3[2]
450 ); 523 );
451 - glVertex2f(v_x+len, v_y+len); 524 + glVertex2f(v_x+t_length, v_y+t_length);
452 glTexCoord3f( 525 glTexCoord3f(
453 p4[0], 526 p4[0],
454 p4[1], 527 p4[1],
455 p4[2] 528 p4[2]
456 ); 529 );
457 - glVertex2f(v_x, v_y+len); 530 + glVertex2f(v_x, v_y+t_length);
458 glEnd(); 531 glEnd();
459 532
460 p1 = ver.p(1,1); 533 p1 = ver.p(1,1);
@@ -467,25 +540,25 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -467,25 +540,25 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
467 p1[1], 540 p1[1],
468 p1[2] 541 p1[2]
469 ); 542 );
470 - glVertex2f(v_x+len, v_y); 543 + glVertex2f(v_x+t_length, v_y);
471 glTexCoord3f( 544 glTexCoord3f(
472 p2[0], 545 p2[0],
473 p2[1], 546 p2[1],
474 p2[2] 547 p2[2]
475 ); 548 );
476 - glVertex2f(v_x+2.0*len, v_y); 549 + glVertex2f(v_x+2.0*t_length, v_y);
477 glTexCoord3f( 550 glTexCoord3f(
478 p3[0], 551 p3[0],
479 p3[1], 552 p3[1],
480 p3[2] 553 p3[2]
481 ); 554 );
482 - glVertex2f(v_x+2.0*len, v_y+len); 555 + glVertex2f(v_x+2.0*t_length, v_y+t_length);
483 glTexCoord3f( 556 glTexCoord3f(
484 p4[0], 557 p4[0],
485 p4[1], 558 p4[1],
486 p4[2] 559 p4[2]
487 ); 560 );
488 - glVertex2f(v_x+len, v_y+len); 561 + glVertex2f(v_x+t_length, v_y+t_length);
489 glEnd(); 562 glEnd();
490 } 563 }
491 564
@@ -522,7 +595,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -522,7 +595,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
522 delete[] texels; 595 delete[] texels;
523 glBindFramebuffer(GL_FRAMEBUFFER, 0); 596 glBindFramebuffer(GL_FRAMEBUFFER, 0);
524 glBindTexture(GL_TEXTURE_2D, 0); 597 glBindTexture(GL_TEXTURE_2D, 0);
525 - CHECK_OPENGL_ERROR  
526 } 598 }
527 599
528 ///@param uint width sets the width of the buffer. 600 ///@param uint width sets the width of the buffer.
@@ -604,7 +676,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -604,7 +676,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
604 ///Method for controling the buffer and texture binding. 676 ///Method for controling the buffer and texture binding.
605 ///Clears the buffer upon binding. 677 ///Clears the buffer upon binding.
606 void 678 void
607 - Bind(float len = 8.0) 679 + Bind()
608 { 680 {
609 glBindFramebuffer(GL_FRAMEBUFFER, fboID);//set up GL buffer 681 glBindFramebuffer(GL_FRAMEBUFFER, fboID);//set up GL buffer
610 glFramebufferTexture2D( 682 glFramebufferTexture2D(
@@ -623,8 +695,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -623,8 +695,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
623 glLoadIdentity(); 695 glLoadIdentity();
624 glMatrixMode(GL_MODELVIEW); 696 glMatrixMode(GL_MODELVIEW);
625 glLoadIdentity(); 697 glLoadIdentity();
626 - glViewport(0,0,2.0*len, numSamples*len);  
627 - gluOrtho2D(0.0,2.0*len,0.0,numSamples*len); 698 + glViewport(0,0,2.0*t_length, numSamples*t_length);
  699 + gluOrtho2D(0.0,2.0*t_length,0.0,numSamples*t_length);
628 glEnable(GL_TEXTURE_3D); 700 glEnable(GL_TEXTURE_3D);
629 glBindTexture(GL_TEXTURE_3D, texID); 701 glBindTexture(GL_TEXTURE_3D, texID);
630 702
@@ -639,6 +711,11 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -639,6 +711,11 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
639 void 711 void
640 Bind(GLuint &textureID, GLuint &framebufferID, int nSamples, float len = 8.0) 712 Bind(GLuint &textureID, GLuint &framebufferID, int nSamples, float len = 8.0)
641 { 713 {
  714 +
  715 +// std::cout << glGetIntegerv(GL_MAX_TEXTURE_SIZE) << std::endl;
  716 + GLint max;
  717 + glGetIntegerv(GL_MAX_TEXTURE_SIZE, &max);
  718 +// std::cout << max << std::endl;
642 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID);//set up GL buffer 719 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID);//set up GL buffer
643 glFramebufferTexture2D( 720 glFramebufferTexture2D(
644 GL_FRAMEBUFFER, 721 GL_FRAMEBUFFER,
@@ -742,7 +819,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -742,7 +819,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
742 start = std::clock(); 819 start = std::clock();
743 #endif 820 #endif
744 stim::vec<int> cost = 821 stim::vec<int> cost =
745 - stim::cuda::get_cost(tID, GL_TEXTURE_2D, n); 822 + stim::cuda::get_cost(tID, GL_TEXTURE_2D, n, 2*t_length, t_length);
746 cudaDeviceSynchronize(); 823 cudaDeviceSynchronize();
747 #ifdef TESTING 824 #ifdef TESTING
748 duration_cuda = duration_cuda + 825 duration_cuda = duration_cuda +
@@ -760,6 +837,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -760,6 +837,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
760 initCuda() 837 initCuda()
761 { 838 {
762 stim::cudaSetDevice(); 839 stim::cudaSetDevice();
  840 + MonteCarloDirectionVectors(500);
763 //GLint max; 841 //GLint max;
764 //glGetIntegerv(GL_MAX_TEXTURE_SIZE, &max); 842 //glGetIntegerv(GL_MAX_TEXTURE_SIZE, &max);
765 //std::cout << max << std::endl; 843 //std::cout << max << std::endl;
@@ -790,12 +868,15 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -790,12 +868,15 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
790 gl_spider 868 gl_spider
791 (int samples = 1089, int samplespos = 441,int samplesmag = 144) 869 (int samples = 1089, int samplespos = 441,int samplesmag = 144)
792 { 870 {
  871 +// std::cout << "I ran this constructor" << std::endl;
793 p = vec<float>(0.0, 0.0, 0.0); 872 p = vec<float>(0.0, 0.0, 0.0);
794 d = vec<float>(0.0, 0.0, 1.0); 873 d = vec<float>(0.0, 0.0, 1.0);
795 m = vec<float>(1.0, 1.0); 874 m = vec<float>(1.0, 1.0);
796 S = vec<float>(1.0, 1.0, 1.0); 875 S = vec<float>(1.0, 1.0, 1.0);
797 R = vec<float>(1.0, 1.0, 1.0); 876 R = vec<float>(1.0, 1.0, 1.0);
  877 +// std::cout << samples << std::endl;
798 numSamples = samples; 878 numSamples = samples;
  879 +// std::cout << numSamples << std::endl;
799 numSamplesPos = samplespos; 880 numSamplesPos = samplespos;
800 numSamplesMag = samplesmag; 881 numSamplesMag = samplesmag;
801 } 882 }
@@ -863,24 +944,29 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -863,24 +944,29 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
863 void 944 void
864 attachSpider(GLuint id) 945 attachSpider(GLuint id)
865 { 946 {
  947 + srand(100);
866 texID = id; 948 texID = id;
867 //GenerateFBO(16, numSamples*8); 949 //GenerateFBO(16, numSamples*8);
868 - GenerateFBO(16, numSamples*8, texbufferID, fboID);  
869 - GenerateFBO(16, numSamplesPos*8, ptexbufferID, pfboID);  
870 - GenerateFBO(16, numSamplesMag*8, mtexbufferID, mfboID); 950 + GenerateFBO(t_length*2, numSamples*t_length, texbufferID, fboID);
  951 + CHECK_OPENGL_ERROR
  952 + GenerateFBO(t_length*2, numSamplesPos*t_length, ptexbufferID, pfboID);
  953 + CHECK_OPENGL_ERROR
  954 + GenerateFBO(t_length*2, numSamplesMag*t_length, mtexbufferID, mfboID);
  955 + CHECK_OPENGL_ERROR
871 GenerateFBO(16, 216, btexbufferID, bfboID); 956 GenerateFBO(16, 216, btexbufferID, bfboID);
  957 + CHECK_OPENGL_ERROR
872 // setDims(0.6, 0.6, 1.0); 958 // setDims(0.6, 0.6, 1.0);
873 // setSize(1024.0, 1024.0, 1024.0); 959 // setSize(1024.0, 1024.0, 1024.0);
874 setMatrix(); 960 setMatrix();
875 dList = glGenLists(3); 961 dList = glGenLists(3);
876 glListBase(dList); 962 glListBase(dList);
877 - Bind(texbufferID, fboID, numSamples); 963 + Bind(texbufferID, fboID, numSamples, t_length);
878 genDirectionVectors(5*M_PI/4); 964 genDirectionVectors(5*M_PI/4);
879 Unbind(); 965 Unbind();
880 - Bind(ptexbufferID, pfboID, numSamplesPos); 966 + Bind(ptexbufferID, pfboID, numSamplesPos, t_length);
881 genPositionVectors(); 967 genPositionVectors();
882 Unbind(); 968 Unbind();
883 - Bind(mtexbufferID, mfboID, numSamplesMag); 969 + Bind(mtexbufferID, mfboID, numSamplesMag, t_length);
884 genMagnitudeVectors(); 970 genMagnitudeVectors();
885 Unbind(); 971 Unbind();
886 Bind(btexbufferID, bfboID, 27); 972 Bind(btexbufferID, bfboID, 27);
@@ -1058,7 +1144,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1058,7 +1144,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1058 void 1144 void
1059 setSeedMag(float mag) 1145 setSeedMag(float mag)
1060 { 1146 {
1061 - std::cout << "sMag: " << mag << std::endl; 1147 +// std::cout << "sMag: " << mag << std::endl;
1062 seedsmags.push(mag); 1148 seedsmags.push(mag);
1063 } 1149 }
1064 1150
@@ -1071,7 +1157,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1071,7 +1157,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1071 void 1157 void
1072 setSeed(float x, float y, float z) 1158 setSeed(float x, float y, float z)
1073 { 1159 {
1074 - std::cout << "sPos: " << x << " " << y << " " << z << std::endl; 1160 +// std::cout << "sPos: " << x << " " << y << " " << z << std::endl;
1075 seeds.push(stim::vec<float>(x, y, z)); 1161 seeds.push(stim::vec<float>(x, y, z));
1076 } 1162 }
1077 1163
@@ -1082,7 +1168,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1082,7 +1168,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1082 void 1168 void
1083 setSeedVec(float x, float y, float z) 1169 setSeedVec(float x, float y, float z)
1084 { 1170 {
1085 - std::cout << "sDir: " << x << " " << y << " " << z << std::endl; 1171 +// std::cout << "sDir: " << x << " " << y << " " << z << std::endl;
1086 seedsvecs.push(stim::vec<float>(x, y, z)); 1172 seedsvecs.push(stim::vec<float>(x, y, z));
1087 } 1173 }
1088 1174
@@ -1222,17 +1308,17 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1222,17 +1308,17 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1222 int 1308 int
1223 Step() 1309 Step()
1224 { 1310 {
1225 - Bind(texbufferID, fboID, numSamples); 1311 + Bind(texbufferID, fboID, numSamples, t_length);
1226 CHECK_OPENGL_ERROR 1312 CHECK_OPENGL_ERROR
1227 #ifdef TESTING 1313 #ifdef TESTING
1228 start = std::clock(); 1314 start = std::clock();
1229 #endif 1315 #endif
1230 findOptimalDirection(); 1316 findOptimalDirection();
1231 Unbind(); 1317 Unbind();
1232 - Bind(ptexbufferID, pfboID, numSamplesPos); 1318 + Bind(ptexbufferID, pfboID, numSamplesPos, t_length);
1233 findOptimalPosition(); 1319 findOptimalPosition();
1234 Unbind(); 1320 Unbind();
1235 - Bind(mtexbufferID, mfboID, numSamplesMag); 1321 + Bind(mtexbufferID, mfboID, numSamplesMag, t_length);
1236 findOptimalScale(); 1322 findOptimalScale();
1237 Unbind(); 1323 Unbind();
1238 CHECK_OPENGL_ERROR 1324 CHECK_OPENGL_ERROR
@@ -1548,6 +1634,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1548,6 +1634,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1548 branchDetection2(); 1634 branchDetection2();
1549 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1635 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1550 addToNetwork(a, spos, smag, sdir); 1636 addToNetwork(a, spos, smag, sdir);
  1637 +// std::cout << "the cost of " << cost << " > " << min_cost << std::endl;
1551 return a; 1638 return a;
1552 break; 1639 break;
1553 } else { 1640 } else {
@@ -1561,6 +1648,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1561,6 +1648,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1561 branchDetection2(); 1648 branchDetection2();
1562 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1649 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1563 addToNetwork(a, spos, smag, sdir); 1650 addToNetwork(a, spos, smag, sdir);
  1651 +// std::cout << "I hit and edge" << std::endl;
1564 return a; 1652 return a;
1565 break; 1653 break;
1566 } 1654 }
@@ -1578,6 +1666,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1578,6 +1666,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1578 branchDetection2(); 1666 branchDetection2();
1579 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1667 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1580 addToNetwork(a, spos, smag, sdir); 1668 addToNetwork(a, spos, smag, sdir);
  1669 +// std::cout << "The templates are too big" << std::endl;
1581 return a; 1670 return a;
1582 break; 1671 break;
1583 } 1672 }
@@ -1586,6 +1675,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1586,6 +1675,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1586 h = selectObject(p, getDirection(), m[0]); 1675 h = selectObject(p, getDirection(), m[0]);
1587 //Have we hit something previously traced? 1676 //Have we hit something previously traced?
1588 if(h != -1){ 1677 if(h != -1){
  1678 +// std::cout << "I hit the fiber " << h << std::endl;
1589 running = false; 1679 running = false;
1590 branchDetection2(); 1680 branchDetection2();
1591 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), h); 1681 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), h);
stim/math/vector.h
@@ -133,6 +133,16 @@ struct vec : public std::vector&lt;T&gt; @@ -133,6 +133,16 @@ struct vec : public std::vector&lt;T&gt;
133 133
134 } 134 }
135 135
  136 +
  137 + vec<T> cyl2cart() const
  138 + {
  139 + vec<T> cyl;
  140 + cyl.push_back(at(0)*std::sin(at(1)));
  141 + cyl.push_back(at(0)*std::cos(at(1)));
  142 + cyl.push_back(at(2));
  143 + return(cyl);
  144 +
  145 + }
136 /// Convert the vector from cartesian to spherical coordinates (x, y, z -> r, theta, phi where theta = [0, 2*pi]) 146 /// Convert the vector from cartesian to spherical coordinates (x, y, z -> r, theta, phi where theta = [0, 2*pi])
137 vec<T> cart2sph() const 147 vec<T> cart2sph() const
138 { 148 {