codebase / stimlib

Commit 0fdb4ed4a7be8cf86aa6f4dffbd412ccdb3e09e7

Authored by Pavel Govyadinov
1 parent 32c433c7

fixed the issue with the improperly visualized textures as well as the cost func… 

…tion. The cost and gl/cuda/gl ping-pong is fully functional
Showing 2 changed files with 22 additions and 10 deletions   Show diff stats
Inline Side-by-side
cuda/cost.h
  Show/Hide comments   View file @0fdb4ed
... ... @@ -46,7 +46,7 @@ float get_sum(float *diff)
46 46  
47 47 __device__ float Template(int x)
48 48 {
49   - if(x < DIM_X/3 || x > DIM_X*2/3){
  49 + if(x < DIM_X/6 || x > DIM_X*5/6 || (x > DIM_X*2/6 && x < DIM_X*4/6)){
50 50 return 1.0;
51 51 }else{
52 52 return 0.0;
... ... @@ -61,18 +61,20 @@ void get_diff (float *result)
61 61 int x = threadIdx.x + blockIdx.x * blockDim.x;
62 62 int y = threadIdx.y + blockIdx.y * blockDim.y;
63 63 int idx = y*DIM_X+x;
  64 + //idx = x*DIM_Y+y;
64 65 //int idx = x*DIM_X+y;
65 66  
66 67 //uchar4 color = tex2D(texIn, x, y);
67 68 //float3 tempcolor = make_float3(color.x, color.y, color.z);
68 69 //float valIn = tempcolor.x + tempcolor.y + tempcolor.z;
69   - float valIn = tex2D(texIn, x, y);
  70 + float valIn = tex2D(texIn, x, y)/255.0;
70 71 float valTemp = Template(x);
71   - result[idx] = valIn;
  72 + result[idx] = abs(valIn-valTemp);
  73 + //-valTemp;
72 74 // - valTemp;
73   - //#if __CUDA_ARCH__>=200
74   - // printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
75   - //#endif
  75 +// #if __CUDA_ARCH__>=200
  76 +// printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
  77 +// #endif
76 78 //cuPrintf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
77 79 }
78 80  
... ...
gl/gl_spider.h
  Show/Hide comments   View file @0fdb4ed
... ... @@ -312,6 +312,8 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
312 312 );
313 313 glVertex2f(1.0, 2.0);
314 314 glEnd();
  315 + //glClearColor(1,1,1,1);
  316 + //glClear(GL_COLOR_BUFFER_BIT);
315 317 glBindTexture(GL_TEXTURE_3D, 0);
316 318 glDisable(GL_TEXTURE_3D);
317 319 glBindFramebuffer(GL_FRAMEBUFFER,0);
... ... @@ -336,6 +338,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
336 338 width, height, 0, GL_LUMINANCE, GL_UNSIGNED_BYTE, texels);
337 339 delete[] texels;
338 340 glBindFramebuffer(GL_FRAMEBUFFER, 0);
  341 + glBindTexture(GL_TEXTURE_2D, 0);
339 342 }
340 343  
341 344  
... ... @@ -353,7 +356,7 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
353 356 printf(" device name: %s\n", prop.name);
354 357 printf(" device maxTexture3D: %d %d %d\n", prop.maxTexture3D[0] ,prop.maxTexture3D[0], prop.maxTexture3D[2]) ;
355 358 HANDLE_ERROR( cudaGLSetGLDevice (device)); */
356   - stim::cudaSetDevice();
  359 + //stim::cudaSetDevice();
357 360 /* HANDLE_ERROR(
358 361 cudaGraphicsMapResources(1, &resource, 0)
359 362 );
... ... @@ -384,10 +387,14 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
384 387 HANDLE_ERROR(
385 388 cudaGraphicsGLRegisterImage(
386 389 &resource,
387   - fboID,
  390 + texbufferID,
388 391 GL_TEXTURE_2D,
389   - CU_GRAPHICS_REGISTER_FLAGS_NONE)
  392 + //CU_GRAPHICS_REGISTER_FLAGS_NONE)
  393 + cudaGraphicsMapFlagsReadOnly)
390 394 );
  395 + //HANDLE_ERROR(
  396 + // cudaGraphicsMapResources(1, &resource)
  397 + //);
391 398 }
392 399  
393 400 void
... ... @@ -396,11 +403,14 @@ class gl_spider : public virtual gl_texture&lt;T&gt;
396 403 HANDLE_ERROR(
397 404 cudaGraphicsUnregisterResource(resource)
398 405 );
  406 + //HANDLE_ERROR(
  407 + // cudaGraphicsUnmapResources(1,&resource)
  408 + //);
399 409 }
400 410  
401 411 float
402 412 getCost()
403   - {
  413 + {
404 414 createResource();
405 415 float cost = get_cost(resource);
406 416 destroyResource();
... ...