Commit 1a456186c7c5524468045633d90e2b32063b08d3

Authored by Pavel Govyadinov
1 parent 0fdb4ed4

Added directional sampling around the spider's direction. All placed into the same buffer

Showing 2 changed files with 178 additions and 53 deletions   Show diff stats
@@ -61,17 +61,12 @@ void get_diff (float *result) @@ -61,17 +61,12 @@ void get_diff (float *result)
61 int x = threadIdx.x + blockIdx.x * blockDim.x; 61 int x = threadIdx.x + blockIdx.x * blockDim.x;
62 int y = threadIdx.y + blockIdx.y * blockDim.y; 62 int y = threadIdx.y + blockIdx.y * blockDim.y;
63 int idx = y*DIM_X+x; 63 int idx = y*DIM_X+x;
64 - //idx = x*DIM_Y+y;  
65 - //int idx = x*DIM_X+y;  
66 64
67 - //uchar4 color = tex2D(texIn, x, y);  
68 - //float3 tempcolor = make_float3(color.x, color.y, color.z);  
69 - //float valIn = tempcolor.x + tempcolor.y + tempcolor.z;  
70 - float valIn = tex2D(texIn, x, y)/255.0; 65 + float valIn = tex2D(texIn, x, y);
  66 + //float valIn = tex2D(texIn, x, y)/255.0;
71 float valTemp = Template(x); 67 float valTemp = Template(x);
72 - result[idx] = abs(valIn-valTemp);  
73 - //-valTemp;  
74 -// - valTemp; 68 + //result[idx] = abs(valIn-valTemp);
  69 + result[idx] = abs(valIn);
75 // #if __CUDA_ARCH__>=200 70 // #if __CUDA_ARCH__>=200
76 // printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]); 71 // printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
77 // #endif 72 // #endif
@@ -6,6 +6,7 @@ @@ -6,6 +6,7 @@
6 #include <cuda.h> 6 #include <cuda.h>
7 #include <cuda_gl_interop.h> 7 #include <cuda_gl_interop.h>
8 #include <cudaGL.h> 8 #include <cudaGL.h>
  9 +#include <math.h>
9 #include "gl_texture.h" 10 #include "gl_texture.h"
10 #include "../visualization/camera.h" 11 #include "../visualization/camera.h"
11 #include "./error.h" 12 #include "./error.h"
@@ -130,8 +131,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -130,8 +131,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
130 attachSpider(GLuint id) 131 attachSpider(GLuint id)
131 { 132 {
132 texID = id; 133 texID = id;
133 - GenerateFBO(400,200);  
134 - Update(); 134 + GenerateFBO(800,400);
  135 + // Update();
  136 + generateVectorField(direction, 75.0);
135 } 137 }
136 138
137 void 139 void
@@ -146,7 +148,42 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -146,7 +148,42 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
146 hor.n()); 148 hor.n());
147 UpdateBuffer(); 149 UpdateBuffer();
148 } 150 }
  151 +
  152 + void
  153 + Update(float v_x, float v_y, vec<float> dir, float dim)
  154 + {
  155 + vec<float> Y(1.0,0.0,0.0);
  156 + if(cos(Y.dot(dir))< 0.087){
  157 + Y[0] = 0.0; Y[1] = 1.0;}
  158 + hor = stim::rect<float>(magnitude, position, dir.norm(),
  159 + ((Y.cross(dir)).cross(dir)).norm());
  160 + ver = stim::rect<float>(magnitude, position, dir.norm(),
  161 + hor.n());
  162 + UpdateBuffer(v_x, v_y, dim);
  163 + }
149 164
  165 + stim::vec<float>
  166 + toSpherical(vec<float> in)
  167 + {
  168 + vec<float> out;
  169 + out[0] = sqrt(pow(in[0],2) + pow(in[1],2) + pow(in[2],2));
  170 + out[1] = atan(in[1]/in[0]);
  171 + out[2] = acos(in[2]/out[0]);
  172 + return out;
  173 + }
  174 +
  175 + stim::vec<float>
  176 + toCartesian(vec<float> in)
  177 + {
  178 + vec<float> out;
  179 + out[0] = cos(in[1])*sin(in[2]);
  180 + out[1] = sin(in[1])*sin(in[2]);
  181 + out[2] = cos(in[2]);
  182 + return out;
  183 + }
  184 +
  185 +
  186 +
150 vec<float> 187 vec<float>
151 getPosition() 188 getPosition()
152 { 189 {
@@ -312,14 +349,145 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -312,14 +349,145 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
312 ); 349 );
313 glVertex2f(1.0, 2.0); 350 glVertex2f(1.0, 2.0);
314 glEnd(); 351 glEnd();
315 - //glClearColor(1,1,1,1); 352 + glBindTexture(GL_TEXTURE_3D, 0);
  353 + glDisable(GL_TEXTURE_3D);
  354 + glBindFramebuffer(GL_FRAMEBUFFER,0);
  355 + glBindTexture(GL_TEXTURE_2D, 0);
  356 + }
  357 +
  358 + void
  359 + UpdateBuffer(float v_x, float v_y, float len)
  360 + {
  361 + std::cout << v_x << "," << v_y << std::endl;
  362 + //float len = dim;
  363 + stim::vec<float>p1;
  364 + stim::vec<float>p2;
  365 + stim::vec<float>p3;
  366 + stim::vec<float>p4;
  367 + glBindFramebuffer(GL_FRAMEBUFFER, fboID);
  368 + glFramebufferTexture2D(
  369 + GL_FRAMEBUFFER,
  370 + GL_COLOR_ATTACHMENT0,
  371 + GL_TEXTURE_2D,
  372 + texbufferID,
  373 + 0);
  374 + glBindFramebuffer(GL_FRAMEBUFFER, fboID);
  375 + GLenum DrawBuffers[1] = {GL_COLOR_ATTACHMENT0};
  376 + glDrawBuffers(1, DrawBuffers);
  377 + glBindTexture(GL_TEXTURE_2D, texbufferID);
  378 + //glClearColor(0,0,0,0);
316 //glClear(GL_COLOR_BUFFER_BIT); 379 //glClear(GL_COLOR_BUFFER_BIT);
  380 + glMatrixMode(GL_PROJECTION);
  381 + glLoadIdentity();
  382 + glMatrixMode(GL_MODELVIEW);
  383 + glLoadIdentity();
  384 + glViewport(0,0,800,400);
  385 + gluOrtho2D(0.0,8.0,0.0,4.0);
  386 + glEnable(GL_TEXTURE_3D);
  387 + glBindTexture(GL_TEXTURE_3D, texID);
  388 + p1 = hor.p(1,1);
  389 + p2 = hor.p(1,0);
  390 + p3 = hor.p(0,0);
  391 + p4 = hor.p(0,1);
  392 + glBegin(GL_QUADS);
  393 + glTexCoord3f(
  394 + p1[0],
  395 + p1[1],
  396 + p1[2]
  397 + );
  398 + //glVertex2f(0.0,0.0);
  399 + glVertex2f(v_x,v_y);
  400 + glTexCoord3f(
  401 + p2[0],
  402 + p2[1],
  403 + p2[2]
  404 + );
  405 + //glVertex2f(1.0, 0.0);
  406 + glVertex2f(v_x+len, v_y);
  407 + glTexCoord3f(
  408 + p3[0],
  409 + p3[1],
  410 + p3[2]
  411 + );
  412 + //glVertex2f(1.0, 2.0);
  413 + glVertex2f(v_x+len, v_y+len);
  414 + glTexCoord3f(
  415 + p4[0],
  416 + p4[1],
  417 + p4[2]
  418 + );
  419 + //glVertex2f(0.0, 2.0);
  420 + glVertex2f(v_x, v_y+len);
  421 + glEnd();
  422 + p1 = ver.p(1,1);
  423 + p2 = ver.p(1,0);
  424 + p3 = ver.p(0,0);
  425 + p4 = ver.p(0,1);
  426 + glBegin(GL_QUADS);
  427 + glTexCoord3f(
  428 + p1[0],
  429 + p1[1],
  430 + p1[2]
  431 + );
  432 + //glVertex2f(1.0, 0.0);
  433 + glVertex2f(v_x+len, v_y);
  434 + glTexCoord3f(
  435 + p2[0],
  436 + p2[1],
  437 + p2[2]
  438 + );
  439 + //glVertex2f(2.0, 0.0);
  440 + glVertex2f(v_x+2*len, v_y);
  441 + glTexCoord3f(
  442 + p3[0],
  443 + p3[1],
  444 + p3[2]
  445 + );
  446 + //glVertex2f(2.0, 2.0);
  447 + glVertex2f(v_x+2*len, v_y+len);
  448 + glTexCoord3f(
  449 + p4[0],
  450 + p4[1],
  451 + p4[2]
  452 + );
  453 + //glVertex2f(1.0, 2.0);
  454 + glVertex2f(v_x+len, v_y+len);
  455 + glEnd();
317 glBindTexture(GL_TEXTURE_3D, 0); 456 glBindTexture(GL_TEXTURE_3D, 0);
318 glDisable(GL_TEXTURE_3D); 457 glDisable(GL_TEXTURE_3D);
319 glBindFramebuffer(GL_FRAMEBUFFER,0); 458 glBindFramebuffer(GL_FRAMEBUFFER,0);
320 glBindTexture(GL_TEXTURE_2D, 0); 459 glBindTexture(GL_TEXTURE_2D, 0);
321 } 460 }
322 - 461 +
  462 + void
  463 + generateVectorField(stim::vec<float> d, float dim)
  464 + {
  465 + vec<float> d_s = toSpherical(d);
  466 + vec<float> temp;
  467 + float Dim = (float) dim;
  468 + float y_0 = 0.0;
  469 + float x_0 = 0.0;
  470 + float len = 4.0/(2.0*Dim+1.0);
  471 + float t0 = M_PI/2;
  472 + float p0 = M_PI/3;
  473 + float dt = 2.0*t0/Dim;
  474 + float dp = 2.0*p0/Dim;
  475 + for(int i = -dim; i <= dim; i++){
  476 + for(int j = -dim; j <= dim; j++){
  477 + //field[i+dim][j+dim][0] = d[0];
  478 + //field[i+dim][j+dim][1] = d[1]+dt*i;
  479 + //field[i+dim][j+dim][2] = d[2]+dp*j;
  480 + temp[0] = 1;
  481 + temp[1] = d_s[1]+dt*i;
  482 + temp[2] = d_s[2]+dp*j;
  483 + temp = toCartesian(temp);
  484 + std::cout << i <<"," << j << std::endl;
  485 + Update(x_0+2.0*(i+dim)*len, y_0+(j+dim)*len, temp, len);
  486 + }
  487 + }
  488 +
  489 + }
  490 +
323 491
324 void 492 void
325 GenerateFBO(unsigned int width, unsigned int height) 493 GenerateFBO(unsigned int width, unsigned int height)
@@ -345,40 +513,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -345,40 +513,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
345 void 513 void
346 initCuda() 514 initCuda()
347 { 515 {
348 - /* cudaDeviceProp prop;  
349 - int device;  
350 - memset( &prop, 0, sizeof(cudaDeviceProp) );  
351 - prop.major = 1;  
352 - prop.minor = 0;  
353 - HANDLE_ERROR( cudaChooseDevice (&device, &prop ) );  
354 - HANDLE_ERROR( cudaGetDeviceProperties(&prop, device));  
355 - printf(" device number: %d\n", device);  
356 - printf(" device name: %s\n", prop.name);  
357 - printf(" device maxTexture3D: %d %d %d\n", prop.maxTexture3D[0] ,prop.maxTexture3D[0], prop.maxTexture3D[2]) ;  
358 - HANDLE_ERROR( cudaGLSetGLDevice (device)); */  
359 - //stim::cudaSetDevice();  
360 -/* HANDLE_ERROR(  
361 - cudaGraphicsMapResources(1, &resource, 0)  
362 - );  
363 - HANDLE_ERROR(  
364 - cudaGraphicsResourceGetMappedPointer(  
365 - &devPtr,  
366 - size,  
367 - resource));  
368 - HANDLE_ERROR(  
369 - cudaGraphicsSubResourceGetMappedArray(  
370 - &c_Array,  
371 - resource,  
372 - 0,0)  
373 - );  
374 - HANDLE_ERROR(  
375 - cudaBindTextureToArray(fboID, c_Array)  
376 - );  
377 - HANDLE_ERROR(  
378 - cudaGraphicsUnmapResources(1, &resource, 0)  
379 - );  
380 - //need to move the constants to video memory.  
381 -*/ 516 + stim::cudaSetDevice();
382 } 517 }
383 518
384 void 519 void
@@ -392,9 +527,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -392,9 +527,6 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
392 //CU_GRAPHICS_REGISTER_FLAGS_NONE) 527 //CU_GRAPHICS_REGISTER_FLAGS_NONE)
393 cudaGraphicsMapFlagsReadOnly) 528 cudaGraphicsMapFlagsReadOnly)
394 ); 529 );
395 - //HANDLE_ERROR(  
396 - // cudaGraphicsMapResources(1, &resource)  
397 - //);  
398 } 530 }
399 531
400 void 532 void
@@ -403,11 +535,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt; @@ -403,11 +535,9 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
403 HANDLE_ERROR( 535 HANDLE_ERROR(
404 cudaGraphicsUnregisterResource(resource) 536 cudaGraphicsUnregisterResource(resource)
405 ); 537 );
406 - //HANDLE_ERROR(  
407 - // cudaGraphicsUnmapResources(1,&resource)  
408 - //);  
409 } 538 }
410 539
  540 +
411 float 541 float
412 getCost() 542 getCost()
413 { 543 {