Commit 556c4e154a958cf575d6563095ad60935d1e3df5

Authored by Pavel Govyadinov
1 parent b710b044

Changed the handling of the transformation matrix in gl_spider.h, added function…

…s for handling the update implementation (compatible with vectors). Added comments and fixes to cost.h, minor changes to matrix.h, and vect.h in order to facilitate moving into 4D vector notation.
Showing 3 changed files with 64 additions and 54 deletions   Show diff stats
stim/cuda/cost.h
... ... @@ -6,13 +6,9 @@
6 6 #include "../visualization/colormap.h"
7 7 #include <sstream>
8 8  
9   -//#define
10   -//#define DIM_Y 10890
11   -//#define DIM_X 20
  9 +
  10 +///Cost function that works with the gl-spider class to find index of the item with min-cost.
12 11 typedef unsigned char uchar;
13   -//surface<void, 2> texOut; ///// maybe just do a normal array instead of a surface.
14   - //we may not need a surface at all.
15   -//texture<float, cudaTextureType2D, cudaReadModeElementType> texTemplate
16 12 texture<uchar, cudaTextureType2D, cudaReadModeElementType> texIn;
17 13 float *result;
18 14 float* v_dif;
... ... @@ -28,7 +24,9 @@ inline void checkCUDAerrors(const char *msg)
28 24 }
29 25 }
30 26  
31   -
  27 +///Finds the sum of all the pixes in a gives template element.
  28 +///Returns the abosolute value.
  29 +///@param *diff, a pointer to the memory block that holds the pixel-differences.
32 30 float get_sum(float *diff)
33 31 {
34 32  
... ... @@ -43,6 +41,9 @@ float get_sum(float *diff)
43 41 return out;
44 42 }
45 43  
  44 +///A virtual representation of a uniform template.
  45 +///Returns the value of the template pixel.
  46 +///@param x, location of a pixel.
46 47 __device__ float Template(int x)
47 48 {
48 49 if(x < 20/6 || x > 20*5/6 || (x > 20*2/6 && x < 20*4/6)){
... ... @@ -53,6 +54,9 @@ __device__ float Template(int x)
53 54  
54 55 }
55 56  
  57 +///Find the difference of the given set of samples and the template
  58 +///using cuda acceleration.
  59 +///@param *result, a pointer to the memory that stores the result.
56 60 __global__
57 61 void get_diff (float *result)
58 62 {
... ... @@ -61,25 +65,17 @@ void get_diff (float *result)
61 65 int y = threadIdx.y + blockIdx.y * blockDim.y;
62 66 int idx = y*20+x;
63 67  
64   - //float valIn = tex2D(texIn, x, y);
65 68 float valIn = tex2D(texIn, x, y)/255.0;
66   - //int idx = x*DIM_X+y;
67   -
68   - //uchar4 color = tex2D(texIn, x, y);
69   - //float3 tempcolor = make_float3(color.x, color.y, color.z);
70   - //float valIn = tempcolor.x + tempcolor.y + tempcolor.z;
71   - //float valIn = tex2D(texIn, x, y);
72 69 float valTemp = Template(x);
73   -// result[idx] = (float) x/(blockDim.x*gridDim.x);//abs(x);
74   - result[idx] = abs(valIn);
75   -// #if __CUDA_ARCH__>=200
76   -// printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
77   -// #endif
  70 + result[idx] = abs(valIn-valTemp);
78 71 }
79 72  
80 73  
81 74  
82   -
  75 +///Initialization function, allocates the memory and passes the necessary
  76 +///handles from OpenGL and Cuda.
  77 +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
  78 +///@param DIM_Y, integer controlling how much memory to allocate.
83 79 void initArray(cudaGraphicsResource_t src, int DIM_Y)
84 80 {
85 81 //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar> ();
... ... @@ -105,7 +101,9 @@ void initArray(cudaGraphicsResource_t src, int DIM_Y)
105 101 // cudaBindTextureToArray(texIn, ptr, &channelDesc)
106 102 // );
107 103 }
108   -
  104 +///Deinit function that frees the memery used and releases the texture resource
  105 +///back to OpenGL.
  106 +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
109 107 void cleanUP(cudaGraphicsResource_t src)
110 108 {
111 109 HANDLE_ERROR(
... ... @@ -121,14 +119,17 @@ void cleanUP(cudaGraphicsResource_t src)
121 119 cudaFree(v_dif)
122 120 );
123 121 }
124   -
  122 +///External access-point to the cuda function
  123 +///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
  124 +///@param DIM_Y, the number of samples in the template.
  125 +///@inter temporary paramenter that tracks the number of times cost.h was called.
125 126 extern "C"
126 127 int get_cost(cudaGraphicsResource_t src, int inter, int DIM_Y)
127 128 {
128 129 float output[DIM_Y];
129 130 float mini = 10000000000000000.0;
130 131 int idx;
131   - stringstream name;
  132 + stringstream name; //for debugging
132 133 initArray(src, DIM_Y*10);
133 134 dim3 grid(20, DIM_Y*10);
134 135 dim3 block(1, 1);
... ...
stim/gl/gl_spider.h
... ... @@ -7,13 +7,14 @@
7 7 #include <cuda_gl_interop.h>
8 8 #include <cudaGL.h>
9 9 #include <math.h>
10   -#include "../gl/gl_texture.h"
11   -#include "../visualization/camera.h"
12   -#include "./error.h"
13   -#include "../math/vector.h"
14   -#include "../math/rect.h"
15   -#include "../cuda/cost.h"
16   -#include "../cuda/glbind.h"
  10 +#include "stim/gl/gl_texture.h"
  11 +#include "stim/visualization/camera.h"
  12 +#include "stim/gl/error.h"
  13 +#include "stim/math/vector.h"
  14 +#include "stim/math/rect.h"
  15 +#include "stim/math/matrix.h"
  16 +#include "stim/cuda/cost.h"
  17 +#include "stim/cuda/glbind.h"
17 18 #include <vector>
18 19  
19 20 #include <iostream>
... ... @@ -42,7 +43,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
42 43 std::vector<stim::vec<float> > dirVectors;
43 44 std::vector<stim::vec<float> > posVectors;
44 45 std::vector<stim::vec<float> > magVectors;
45   - double currentTransform[16];
  46 + stim::matrix<float, 4> currentTransform;
46 47 using gl_texture<T>::texID;
47 48 using gl_texture<T>::S;
48 49 using gl_texture<T>::R;
... ... @@ -62,9 +63,16 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
62 63 {
63 64 genTemplate(dirVectors, 0);
64 65 int best = getCost();
65   - stim::vec<float> next = dirVectors[best];
66   - // next[0] = next[0]*
67   -
  66 + stim::vec<float, 4> next;
  67 + next[0] = dirVectors[best][0]*512.0*0.6;
  68 + next[1] = dirVectors[best][1]*512.0*0.6;
  69 + next[2] = dirVectors[best][2]*512.0*0.6;
  70 + next[3] = 1;
  71 + next = (currentTransform*next).norm();
  72 + setPosition( position[0]+next[0]*magnitude[0]/4,
  73 + position[1]+next[1]*magnitude[0]/4,
  74 + position[2]+next[2]*magnitude[0]/4);
  75 + setDirection(next[0], next[1], next[2]);
68 76 }
69 77  
70 78 /// Method for finding the best direction for the spider.
... ... @@ -419,23 +427,19 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
419 427 stim::vec<float, 4> rot = getRotation(direction);
420 428 glMatrixMode(GL_TEXTURE);
421 429 glLoadIdentity();
422   - glScalef(1.0/512.0/0.6, 1.0/512.0/0.6, 1.0/426.0/2.0);
423   - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform);
424   - printTransform();
  430 + glScalef(1.0/S[1]/R[1], 1.0/S[2]/R[2], 1.0/S[3]/R[3]);
425 431 glTranslatef(position[0],
426 432 position[1],
427 433 position[2]);
428   - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform);
429   - printTransform();
430 434 glScalef(magnitude[0],
431 435 magnitude[1],
432 436 magnitude[0]);
433   - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform);
434   - printTransform();
435 437 glRotatef(rot[0], rot[1], rot[2], rot[3]);
436   - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform);
437   - printTransform();
438   - glGetDoublev(GL_TEXTURE_MATRIX, currentTransform);
  438 + float curTrans[16];
  439 + glGetFloatv(GL_TEXTURE_MATRIX, curTrans);
  440 + fillTransform(curTrans);
  441 + printTransform();
  442 +
439 443 CHECK_OPENGL_ERROR
440 444 glMatrixMode(GL_MODELVIEW);
441 445 }
... ... @@ -499,8 +503,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
499 503 setPosition(0.0,0.0,0.0);
500 504 setDirection(0.0,0.0,1.0);
501 505 setMagnitude(1.0);
502   - //numSamples = 1089;
503   - numSamples = 9;
  506 + numSamples = 1089;
  507 + //numSamples = 9;
504 508 }
505 509 ///@param samples, the number of samples this spider is going to use.
506 510 ///best results if samples is can create a perfect root.
... ... @@ -538,6 +542,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
538 542 genPositionVectors();
539 543 genMagnitudeVectors();
540 544 gl_texture<T>::setDims(0.6, 0.6, 2.0);
  545 + setSize(512, 512, 426);
541 546 genTemplate(dirVectors, 0);
542 547 }
543 548  
... ... @@ -687,15 +692,20 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
687 692 }
688 693  
689 694 void
690   - printTransform()
  695 + fillTransform(float mat[16])
691 696 {
692   - for(int i = 0; i < 4; i++){
693   - std::cout << "[" << currentTransform[i] << "]" <<
694   - "[" << currentTransform[i+4] << "]" <<
695   - "[" << currentTransform[i+8] << "]" <<
696   - "[" << currentTransform[i+12] << "]" <<
697   - std::endl;
  697 + for(int r = 0; r < 4; r++){
  698 + for(int c = 0; c < 4; c++){
  699 + currentTransform(r,c) = mat[c*4+r];
698 700 }
  701 + }
  702 +
  703 + }
  704 +
  705 + void
  706 + printTransform()
  707 + {
  708 + std::cout << currentTransform << std::endl;
699 709 }
700 710  
701 711 /* Method for initializing the cuda devices, necessary only
... ...
stim/math/vector.h
... ... @@ -203,7 +203,6 @@ struct vec
203 203 v[i] = rhs;
204 204 return *this;
205 205 }
206   -
207 206 template<typename Y>
208 207 CUDA_CALLABLE vec<T, N> & operator=(vec<Y, N> rhs){
209 208 for(int i=0; i<N; i++)
... ...