Трассировка лучей на GPU

День добрый,
Решил я тут написать программу трассировка лучей из 5 главы книги "Технология CUDA в примерах"
Вот код(как сделать его в бб-кодах я не ведаю):

  1. <br />
  2. #include "stdafx.h"<br />
  3. #include "cuda.h"<br />
  4. #include "cuda_runtime_api.h"<br />
  5. #include "book.h"<br />
  6. #include "cpu_bitmap.h"<br />
  7. #include <stdlib.h><br />
  8. #include <math.h></p>
  9. <p>#define INF 2e10f<br />
  10. #define rnd( x ) (x*rand()/RAND_MAX)<br />
  11. #define SPHERES 20<br />
  12. #define DIM 950</p>
  13. <p>struct Sphere<br />
  14. {<br />
  15.         float r,b,g;<br />
  16.         float radius;<br />
  17.         float x,y,z;</p>
  18. <p>     __device__ float hit( float ox, float oy, float *n)<br />
  19.                 {<br />
  20.                         float dx = ox - x;<br />
  21.                         float dy = oy - y;<br />
  22.                         if(dx*dx + dy*dy < radius*radius)<br />
  23.                                 {<br />
  24.                                         float dz = sqrtf( radius*radius - dx*dx - dy*dy );<br />
  25.                                         *n = dz/sqrtf( radius*radius );<br />
  26.                                         return dz + z;<br />
  27.                                 }<br />
  28.                         return -INF;<br />
  29.                 }<br />
  30. };</p>
  31. <p>__constant__ Sphere s[SPHERES];<br />
  32. //Sphere s[SPHERES];</p>
  33. <p>__global__ void kernel( unsigned char *ptr )<br />
  34. {<br />
  35.         int x = threadIdx.x + blockIdx.x*blockDim.x;<br />
  36.         int y = threadIdx.y + blockIdx.y*blockDim.y;<br />
  37.         int offset = x + y*blockDim.x*gridDim.x;<br />
  38.         float ox = (x - DIM/2);<br />
  39.         float oy = (y - DIM/2);</p>
  40. <p>     float r=0, g=0, b=0;<br />
  41.         float maxz = -INF;<br />
  42.         for(int i=0; i<SPHERES; i++)<br />
  43.                 {<br />
  44.                         float n;<br />
  45.                         float t = s[i].hit( ox, oy, &n );<br />
  46.                         if(t > maxz)<br />
  47.                                 {<br />
  48.                                         float fscale = n;<br />
  49.                                         r = s[i].r * fscale;<br />
  50.                                         g = s[i].g * fscale;<br />
  51.                                         b = s[i].b * fscale;<br />
  52.                                         maxz = t;<br />
  53.                                 }<br />
  54.                 }</p>
  55. <p>     ptr[offset*4 + 0] = (int)(r*255);<br />
  56.         ptr[offset*4 + 1] = (int)(g*255);<br />
  57.         ptr[offset*4 + 2] = (int)(b*255);<br />
  58.         ptr[offset*4 + 3] = 255;<br />
  59. }</p>
  60. <p>int _tmain(int argc, _TCHAR* argv[])<br />
  61. {<br />
  62.         //запомнить время начала<br />
  63.         cudaEvent_t start,stop;<br />
  64.         HANDLE_ERROR( cudaEventCreate( &start ) );<br />
  65.         HANDLE_ERROR( cudaEventCreate( &stop ) );<br />
  66.         HANDLE_ERROR( cudaEventRecord( start, 0 ) );</p>
  67. <p>     CPUBitmap bitmap( DIM, DIM );<br />
  68.         unsigned char *dev_bitmap;</p>
  69. <p>     //выделить на GPU память для выходного изображения<br />
  70.         HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap , bitmap.image_size() ) );</p>
  71. <p>     //выделить память для набора схем<br />
  72.         //HANDLE_ERROR( cudaMalloc( (void**)&s, sizeof(Sphere)*SPHERES ) );</p>
  73. <p>     //выделить временную память,инициализировать её, скопировать на GPU, а потом освободить<br />
  74.         Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere)*SPHERES );<br />
  75.         for(int i=0; i<SPHERES; i++)<br />
  76.                 {<br />
  77.                         temp_s[i].r = rnd( 1.0f );<br />
  78.                         temp_s[i].b = rnd( 1.0f );<br />
  79.                         temp_s[i].g = rnd( 1.0f );<br />
  80.                         temp_s[i].x = rnd( 1000.0f ) - 500;<br />
  81.                         temp_s[i].y = rnd( 1000.0f ) - 500;<br />
  82.                         temp_s[i].z = rnd( 1000.0f ) - 500;<br />
  83.                         temp_s[i].radius = rnd( 100.0f ) + 20;<br />
  84.                 }</p>
  85. <p>     //HANDLE_ERROR( cudaMemcpy( s, temp_s, sizeof(Sphere)*SPHERES, cudaMemcpyHostToDevice ) );<br />
  86.         HANDLE_ERROR( cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere)*SPHERES ) );<br />
  87.         free( temp_s );</p>
  88. <p>     //сгенерить растровое изображение сфер<br />
  89.         dim3 grids( DIM/16 , DIM/16 );<br />
  90.         dim3 threads(16,16);<br />
  91.         kernel<<<grids,threads>>>( dev_bitmap );</p>
  92. <p>     //скопировать изображение в память CPU для вывода на экран<br />
  93.         HANDLE_ERROR( cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost ) );<br />
  94.         bitmap.display_and_exit();</p>
  95. <p>     cudaFree( dev_bitmap );<br />
  96.         //cudaFree( s );</p>
  97. <p>     return 0;<br />
  98. }<br />

у меня стоит студия2010 + NSight 2.2. TDR через NSight я выключил. Это вроде хорошо уже)
Теперь компилирую этот проект через NSight->Start Graphics debugging
в результате он мне выдает на картинке вместо нужных кругов какие-то наклоненные полосы непонятной формы,что не есть хорошо. А сама программа NSight Monitor выдает такое сообщение:
>>Nsight Debug
Shader Debugging and Pixel History are disabled when running locally.

Что мне делать?
Подскажите пожалуйста и желательно с пошаговой инструкцией

Forums: