Commit e45b97ced627871662615e182099a0dfba38cd7f

Authored by Pavel Govyadinov
1 parent d06d3dbd

safety commit, I broke something and couldn't go back because I didn't commit

stim/cuda/branch_detection.cuh
@@ -111,6 +111,7 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y) @@ -111,6 +111,7 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y)
111 111
112 112
113 113
  114 + test(texbufferID, texType, x, y);
114 std::vector<stim::vec<float> > output; 115 std::vector<stim::vec<float> > output;
115 initCuda(bytes_table, bytes_ds); 116 initCuda(bytes_table, bytes_ds);
116 117
@@ -123,6 +124,8 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y) @@ -123,6 +124,8 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y)
123 stim::cuda::tex_gaussian_blur2<float>( 124 stim::cuda::tex_gaussian_blur2<float>(
124 gpuI, sigma, x, y, t.getTexture(), t.getArray() 125 gpuI, sigma, x, y, t.getTexture(), t.getArray()
125 ); 126 );
  127 + stim::gpu2image<float>(gpuI, "Blur.jpg", x,y , 0, 255);
  128 +// stim::gpu2image<float>(t.getArray(), "ORIGINAL.jpg", x,y , 0, 255);
126 cudaDeviceSynchronize(); 129 cudaDeviceSynchronize();
127 130
128 131
@@ -146,12 +149,14 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y) @@ -146,12 +149,14 @@ find_branch(GLint texbufferID, GLenum texType, unsigned int x, unsigned int y)
146 149
147 cudaDeviceSynchronize(); 150 cudaDeviceSynchronize();
148 stim::cuda::gpu_local_max<float>(gpuCenters, gpuVote, final_t, conn, x, y); 151 stim::cuda::gpu_local_max<float>(gpuCenters, gpuVote, final_t, conn, x, y);
  152 + stim::gpu2image<float>(gpuCenters, "Centers.jpg", x, y, 0, 1);
149 cudaMemcpy(cpuCenters, gpuCenters, bytes_ds, cudaMemcpyDeviceToHost); 153 cudaMemcpy(cpuCenters, gpuCenters, bytes_ds, cudaMemcpyDeviceToHost);
  154 + stim::cpu2image<float>(cpuCenters, "CentersXPU.jpg", x, y, 0, 1);
150 for(int i = 0; i < pixels; i++) 155 for(int i = 0; i < pixels; i++)
151 { 156 {
152 int ix = (i % x); 157 int ix = (i % x);
153 int iy = (i / x); 158 int iy = (i / x);
154 - if((cpuCenters[i] == 1) && (ix > 4) && (ix < x-4)) 159 + if((cpuCenters[i] != 0))
155 { 160 {
156 161
157 float x_v = (float) ix; 162 float x_v = (float) ix;
stim/cuda/testKernel.cuh
@@ -14,9 +14,9 @@ @@ -14,9 +14,9 @@
14 ///Initialization function, allocates the memory and passes the necessary 14 ///Initialization function, allocates the memory and passes the necessary
15 ///handles from OpenGL and Cuda. 15 ///handles from OpenGL and Cuda.
16 ///@param DIM_Y --integer controlling how much memory to allocate. 16 ///@param DIM_Y --integer controlling how much memory to allocate.
17 - void initArray() 17 + void initArray(int x, int y)
18 { 18 {
19 - cudaMalloc( (void**) &print, 216*16*sizeof(float)); ///temporary 19 + cudaMalloc( (void**) &print, x*y*sizeof(float)); ///temporary
20 } 20 }
21 21
22 ///Deinit function that frees the memery used and releases the texture resource 22 ///Deinit function that frees the memery used and releases the texture resource
@@ -48,12 +48,14 @@ @@ -48,12 +48,14 @@
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*16+x; 51 + int idx = y*64+x;
  52 +// int idx = y*32+x;
  53 +// int idx = y*16+x;
52 54
53 float valIn = tex2D<unsigned char>(texIn, x, y); 55 float valIn = tex2D<unsigned char>(texIn, x, y);
54 float templa = templ(x); 56 float templa = templ(x);
55 - //print[idx] = abs(valIn); ///temporary  
56 - print[idx] = abs(templa); ///temporary 57 + print[idx] = valIn; ///temporary
  58 + //print[idx] = abs(templa); ///temporary
57 59
58 } 60 }
59 61
@@ -64,7 +66,7 @@ @@ -64,7 +66,7 @@
64 ///@param GLenum texType --either GL_TEXTURE_1D, GL_TEXTURE_2D or GL_TEXTURE_3D 66 ///@param GLenum texType --either GL_TEXTURE_1D, GL_TEXTURE_2D or GL_TEXTURE_3D
65 /// may work with other gl texture types, but untested. 67 /// may work with other gl texture types, but untested.
66 ///@param DIM_Y, the number of samples in the template. 68 ///@param DIM_Y, the number of samples in the template.
67 - void test(GLint texbufferID, GLenum texType) 69 + void test(GLint texbufferID, GLenum texType, int x, int y)
68 { 70 {
69 71
70 //Bind the Texture in GL and allow access to cuda. 72 //Bind the Texture in GL and allow access to cuda.
@@ -72,16 +74,12 @@ @@ -72,16 +74,12 @@
72 74
73 //initialize the return arrays. 75 //initialize the return arrays.
74 76
75 - initArray();  
76 -  
77 - int x = 16;  
78 - int y = 27*8;  
79 - y = 8* 1089; 77 + initArray(x,y);
  78 + dim3 numBlocks(1, y);
  79 + dim3 threadsPerBlock(x, 1);
80 int max_threads = stim::maxThreadsPerBlock(); 80 int max_threads = stim::maxThreadsPerBlock();
81 //dim3 threads(max_threads, 1); 81 //dim3 threads(max_threads, 1);
82 //dim3 blocks(x / threads.x + 1, y); 82 //dim3 blocks(x / threads.x + 1, y);
83 - dim3 numBlocks(1, 1089);  
84 - dim3 threadsPerBlock(16, 8);  
85 //dim3 numBlocks(2, 2); 83 //dim3 numBlocks(2, 2);
86 //dim3 threadsPerBlock(8, 108); 84 //dim3 threadsPerBlock(8, 108);
87 85
@@ -92,7 +90,7 @@ @@ -92,7 +90,7 @@
92 cudaDeviceSynchronize(); 90 cudaDeviceSynchronize();
93 stringstream name; //for debugging 91 stringstream name; //for debugging
94 name << "FromTex.bmp"; 92 name << "FromTex.bmp";
95 - stim::gpu2image<float>(print, name.str(),16,1089*8,0,1.0); 93 + stim::gpu2image<float>(print, name.str(),x,y,0,255);
96 94
97 tx.UnmapCudaTexture(); 95 tx.UnmapCudaTexture();
98 cleanUP(); 96 cleanUP();
stim/gl/gl_spider.h
@@ -23,6 +23,7 @@ @@ -23,6 +23,7 @@
23 #include <stim/cuda/branch_detection.cuh> 23 #include <stim/cuda/branch_detection.cuh>
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/cuda/testKernel.cuh> 27 //#include <stim/cuda/testKernel.cuh>
27 28
28 //#include <stim/cuda/testKernel.cuh> 29 //#include <stim/cuda/testKernel.cuh>
@@ -218,6 +219,37 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -218,6 +219,37 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
218 219
219 } 220 }
220 221
  222 +
  223 + void
  224 + branchDetection2(int n = 8, int l_template = 8, int l_square = 8)
  225 + {
  226 + if(cL.size() < 4){}
  227 + else{
  228 + setMatrix(1);
  229 + DrawLongCylinder(n, l_template, l_square);
  230 + stim::cylinder<float> cyl(cL, cM);
  231 + std::vector< stim::vec<float> > result = find_branch(btexbufferID, GL_TEXTURE_2D, n*l_square, (cL.size()-1)*l_template);
  232 + std::cerr << "the number of points is " << result.size() << std::endl;
  233 + if(!result.empty())
  234 + {
  235 + for(int i = 0; i < result.size(); i++)
  236 + {
  237 + std::cout << "Testing "<< i << ": " << result[i][0] << ", " << result[i][1] << std::endl;
  238 + stim::vec<float> v = cyl.surf(result[i][0], result[i][1]*360);
  239 + std::cout << v[0] << " ," << v[1] << " ," << v[2] << std::endl;
  240 + stim::vec<float> di = cyl.p(result[i][0]);
  241 + std::cout << di[0] << " ," << di[1] << " ," << di[2] << std::endl;
  242 + float rad = cyl.r(result[i][0]);
  243 + std::cout << rad << std::endl;
  244 + setSeed(v);
  245 + setSeedVec(stim::vec<float>(-di[0]+v[0], -di[1]+v[1], -di[2]+v[2]).norm());
  246 + setSeedMag(cyl.r(result[i][0]));
  247 + }
  248 + }
  249 + std::cout << "I ran the new branch detection" << std::endl;
  250 + }
  251 + }
  252 +
221 253
222 //--------------------------------------------------------------------------// 254 //--------------------------------------------------------------------------//
223 //---------------------TEMPLATE CREATION METHODS----------------------------// 255 //---------------------TEMPLATE CREATION METHODS----------------------------//
@@ -459,6 +491,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -459,6 +491,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
459 void 491 void
460 GenerateFBO(unsigned int width, unsigned int height, GLuint &textureID, GLuint &framebufferID) 492 GenerateFBO(unsigned int width, unsigned int height, GLuint &textureID, GLuint &framebufferID)
461 { 493 {
  494 + glDeleteFramebuffers(1, &framebufferID);
462 glGenFramebuffers(1, &framebufferID); 495 glGenFramebuffers(1, &framebufferID);
463 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID); 496 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID);
464 int numChannels = 1; 497 int numChannels = 1;
@@ -506,45 +539,60 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -506,45 +539,60 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
506 } 539 }
507 540
508 541
  542 + ///IF type == 0
509 ///Method for using the gl manipulation to align templates from 543 ///Method for using the gl manipulation to align templates from
510 ///Template space (-0.5 0.5) to Texture space (0.0, 1.0), 544 ///Template space (-0.5 0.5) to Texture space (0.0, 1.0),
511 ///Based on the p of the spider in real space (arbitrary). 545 ///Based on the p of the spider in real space (arbitrary).
  546 +
  547 + ///IF type == 1
  548 + ///Method for using the gl manipulation to set up a matrix
  549 + ///To transform from tissue space into texture space.
  550 + ///All transformation happen in glMatrixMode(GL_TEXTURE).
512 ///All transformation happen in glMatrixMode(GL_TEXTURE). 551 ///All transformation happen in glMatrixMode(GL_TEXTURE).
513 - void setMatrix() 552 + void setMatrix(int type = 0)
514 { 553 {
515 - float curTrans[16]; //array to store the matrix values.  
516 - stim::vec<float> rot = getRotation(d); //get the rotation parameters for the current direction vector.  
517 - glMatrixMode(GL_TEXTURE);  
518 - glLoadIdentity(); 554 + if(type == 0)
  555 + {
  556 + float curTrans[16]; //array to store the matrix values.
  557 + stim::vec<float> rot = getRotation(d); //get the rotation parameters for the current direction vector.
  558 + glMatrixMode(GL_TEXTURE);
  559 + glLoadIdentity();
519 560
520 - //Scale by the voxel size and number of slices.  
521 - glScalef(1.0/S[0]/R[0], 1.0/S[1]/R[1], 1.0/S[2]/R[2]);  
522 - //translate to the current position of the spider in the texture.  
523 - glTranslatef(p[0],  
524 - p[1],  
525 - p[2]);  
526 - //rotate to the current direction of the spider.  
527 - glRotatef(rot[0], rot[1], rot[2], rot[3]);  
528 - //scale to the magnitude of the spider.  
529 - glScalef(m[0],  
530 - m[0],  
531 - m[0]);  
532 - //get and store the current transformation matrix for later use.  
533 - glGetFloatv(GL_TEXTURE_MATRIX, curTrans);  
534 - cT.set(curTrans);  
535 - // printTransform();  
536 -  
537 - CHECK_OPENGL_ERROR  
538 - //revert back to default gl mode.  
539 - glMatrixMode(GL_MODELVIEW); 561 + //Scale by the voxel size and number of slices.
  562 + glScalef(1.0/S[0]/R[0], 1.0/S[1]/R[1], 1.0/S[2]/R[2]);
  563 + //translate to the current position of the spider in the texture.
  564 + glTranslatef(p[0],
  565 + p[1],
  566 + p[2]);
  567 + //rotate to the current direction of the spider.
  568 + glRotatef(rot[0], rot[1], rot[2], rot[3]);
  569 + //scale to the magnitude of the spider.
  570 + glScalef(m[0],
  571 + m[0],
  572 + m[0]);
  573 + //get and store the current transformation matrix for later use.
  574 + glGetFloatv(GL_TEXTURE_MATRIX, curTrans);
  575 + cT.set(curTrans);
  576 + // printTransform();
  577 +
  578 + CHECK_OPENGL_ERROR
  579 + //revert back to default gl mode.
  580 + glMatrixMode(GL_MODELVIEW);
  581 + }
  582 + else if(type == 1)
  583 + {
  584 + glMatrixMode(GL_TEXTURE);
  585 + glLoadIdentity();
  586 + glScalef(1.0/S[0]/R[0], 1.0/S[1]/R[1], 1.0/S[2]/R[2]);
  587 + glMatrixMode(GL_MODELVIEW);
  588 + }
540 } 589 }
541 590
542 ///Method for controling the buffer and texture binding. 591 ///Method for controling the buffer and texture binding.
543 ///Clears the buffer upon binding. 592 ///Clears the buffer upon binding.
544 void 593 void
545 - Bind() 594 + Bind(float len = 8.0)
546 { 595 {
547 - float len = 8.0;  
548 glBindFramebuffer(GL_FRAMEBUFFER, fboID);//set up GL buffer 596 glBindFramebuffer(GL_FRAMEBUFFER, fboID);//set up GL buffer
549 glFramebufferTexture2D( 597 glFramebufferTexture2D(
550 GL_FRAMEBUFFER, 598 GL_FRAMEBUFFER,
@@ -576,9 +624,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -576,9 +624,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
576 ///@param GLuint &framebufferID, framebuffer used for storage. 624 ///@param GLuint &framebufferID, framebuffer used for storage.
577 ///@param int nSamples, number of rectanges to create. 625 ///@param int nSamples, number of rectanges to create.
578 void 626 void
579 - Bind(GLuint &textureID, GLuint &framebufferID, int nSamples) 627 + Bind(GLuint &textureID, GLuint &framebufferID, int nSamples, float len = 8.0)
580 { 628 {
581 - float len = 8.0;  
582 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID);//set up GL buffer 629 glBindFramebuffer(GL_FRAMEBUFFER, framebufferID);//set up GL buffer
583 glFramebufferTexture2D( 630 glFramebufferTexture2D(
584 GL_FRAMEBUFFER, 631 GL_FRAMEBUFFER,
@@ -1216,6 +1263,36 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1216,6 +1263,36 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1216 glEnd(); 1263 glEnd();
1217 glEndList(); 1264 glEndList();
1218 } 1265 }
  1266 +
  1267 + void
  1268 + DrawLongCylinder(int n = 8, int l_template = 8,int l_square = 8)
  1269 + {
  1270 + int cylLen = cL.size()-1;
  1271 + GenerateFBO(n*l_square, cylLen*l_template, btexbufferID, bfboID);
  1272 + Bind(btexbufferID, bfboID, cylLen, l_template*l_square/2.0);
  1273 + stim::cylinder<float> cyl(cL, cM);
  1274 + std::vector<std::vector<stim::vec<float> > > p = cyl.getPoints(n);
  1275 + for(int i = 0; i < p.size()-1; i++) ///number of circles
  1276 + {
  1277 + for(int j = 0; j < p[0].size()-1; j++) ///points in the circle
  1278 + {
  1279 + glBegin(GL_QUADS);
  1280 + glTexCoord3f(p[i][j][0], p[i][j][1], p[i][j][2]);
  1281 + glVertex2f(j*l_square, i*(float)l_template);
  1282 +
  1283 + glTexCoord3f(p[i][j+1][0], p[i][j+1][1], p[i][j+1][2]);
  1284 + glVertex2f(j*l_square+l_square, i*(float)l_template);
  1285 +
  1286 + glTexCoord3f(p[i+1][j+1][0], p[i+1][j+1][1], p[i+1][j+1][2]);
  1287 + glVertex2f(j*l_square+l_square, i*(float)l_template+(float)l_template);
  1288 +
  1289 + glTexCoord3f(p[i+1][j][0], p[i+1][j][1], p[i+1][j][2]);
  1290 + glVertex2f(j*l_square,i*(float)l_template+(float)l_template);
  1291 + glEnd();
  1292 + }
  1293 + }
  1294 + Unbind();
  1295 + }
1219 1296
1220 1297
1221 ///@param min_cost the cost value used for tracing 1298 ///@param min_cost the cost value used for tracing
@@ -1515,6 +1592,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1515,6 +1592,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1515 if (cost > min_cost){ 1592 if (cost > min_cost){
1516 running = false; 1593 running = false;
1517 sk.End(); 1594 sk.End();
  1595 + branchDetection2();
1518 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1596 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1519 addToNetwork(a, spos, smag, sdir); 1597 addToNetwork(a, spos, smag, sdir);
1520 return a; 1598 return a;
@@ -1529,6 +1607,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1529,6 +1607,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1529 // std::cout << "Found Edge" << std::endl; 1607 // std::cout << "Found Edge" << std::endl;
1530 running = false; 1608 running = false;
1531 sk.End(); 1609 sk.End();
  1610 + branchDetection2();
1532 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1611 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1533 addToNetwork(a, spos, smag, sdir); 1612 addToNetwork(a, spos, smag, sdir);
1534 return a; 1613 return a;
@@ -1548,6 +1627,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1548,6 +1627,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1548 // std::cout << "Magnitude Limit" << std::endl; 1627 // std::cout << "Magnitude Limit" << std::endl;
1549 running = false; 1628 running = false;
1550 sk.End(); 1629 sk.End();
  1630 + branchDetection2();
1551 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1); 1631 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), -1);
1552 addToNetwork(a, spos, smag, sdir); 1632 addToNetwork(a, spos, smag, sdir);
1553 return a; 1633 return a;
@@ -1560,6 +1640,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1560,6 +1640,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1560 if(h != -1){ 1640 if(h != -1){
1561 running = false; 1641 running = false;
1562 sk.End(); 1642 sk.End();
  1643 + branchDetection2();
1563 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), h); 1644 pair<stim::fiber<float>, int> a(stim::fiber<float> (cL, cM), h);
1564 addToNetwork(a, spos, smag, sdir); 1645 addToNetwork(a, spos, smag, sdir);
1565 return a; 1646 return a;
@@ -1574,7 +1655,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -1574,7 +1655,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
1574 sk.Vertex(p[0], p[1], p[2]); 1655 sk.Vertex(p[0], p[1], p[2]);
1575 Bind(btexbufferID, bfboID, 27); 1656 Bind(btexbufferID, bfboID, 27);
1576 CHECK_OPENGL_ERROR 1657 CHECK_OPENGL_ERROR
1577 - branchDetection(); 1658 +// branchDetection();
1578 CHECK_OPENGL_ERROR 1659 CHECK_OPENGL_ERROR
1579 Unbind(); 1660 Unbind();
1580 CHECK_OPENGL_ERROR 1661 CHECK_OPENGL_ERROR