From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/RAY/rayTracing.cu | 441 ++++++++++++++++++++++++++++++++++++++ 1 file changed, 441 insertions(+) create mode 100644 benchmarks/CUDA/RAY/rayTracing.cu (limited to 'benchmarks/CUDA/RAY/rayTracing.cu') diff --git a/benchmarks/CUDA/RAY/rayTracing.cu b/benchmarks/CUDA/RAY/rayTracing.cu new file mode 100644 index 0000000..80145dd --- /dev/null +++ b/benchmarks/CUDA/RAY/rayTracing.cu @@ -0,0 +1,441 @@ +/* + * Copyright 2008 BOROUJERDI Maxime. Tous droits reserves. + */ + +#include +#include +#include +#include +#include "makebmp.h" + +/*#include +#include + +#include */ +#include + +typedef unsigned int uint; +typedef unsigned char uchar; + +#define numObj 4 + +#define PI 3.141592654f +#define Angle(a) ((a*PI)/180.0) + +//#define DEVICE_EMU +//#define DEBUG_RT_CUDA +#define FIXED_CONST_PARSE +#ifdef DEBUG_RT_CUDA +#define DEBUG_NUM 8 +float4 *d_debug_float4; +uint *d_debug_uint; +float4 *h_debug_float4; +uint *h_debug_uint; +#endif +int g_verbose; + +#include + +unsigned width = 64; //640; //512; //16; //32; //512; +unsigned height = 64; //480; //512; //16;//512; +dim3 blockSize(16,8); +dim3 gridSize(width/blockSize.x, height/blockSize.y); + +float3 viewRotation; +float3 viewTranslation = make_float3(0.0, 0.0, -4.0f); +float invViewMatrix[12]; + +//static int fpsCount = 0; // FPS count for averaging +//static int fpsLimit = 1; // FPS limit for sampling +unsigned int timer; + + +//GLuint pbo = 0; // Pixel buffer d'OpenGL + + +void initPixelBuffer(); + +class Observateur +{ + private: + matrice3x4 M; // U, V, W + float df; // distance focale + + public: + Observateur( ); + Observateur(const float3 &, const float3 &, const float3 &, double ); + + inline const matrice3x4 & getMatrice( ) const { return M; } + inline float getDistance( ) const { return df; } +}; + +Observateur::Observateur() +{ + M.m[0] = make_float4(0.0f,0.0f,1.0f,0.0f); + M.m[1] = make_float4(0.0f,1.0f,0.0f,0.0f); + M.m[2] = make_float4(1.0f,0.0f,0.0f,0.0f); + df = 1.0 / tan(Angle(65)/2.0); +} + +Observateur::Observateur(const float3 & p, const float3 & u, const float3 & v, double a ) +{ + float3 VP, U, V, W; + VP = normalize(v); + U = normalize(u); + V = normalize(VP - dot(U,VP)*U); + W = normalize(cross(U,V)); + M.m[0] = make_float4(U.x,U.y,U.z,p.x); + M.m[1] = make_float4(V.x,V.y,V.z,p.y); + M.m[2] = make_float4(W.x,W.y,W.z,p.z); + df = 1.0 / tan(Angle(a)/2.0); +} + +float anim = 0.0f, pas = 0.015f; +Observateur obs = Observateur(make_float3(0.0f,0.5f,2.0f),normalize(make_float3(0.0f,0.0f,0.0f)-make_float3(0.0f,0.5f,2.0f)),make_float3(0.0f,1.0f,0.0f),65.0f);; + +uint * values = NULL, * d_output, * d_temp, NUM; +uint * c_output; + +Node node[numObj], * d_node; + +Sphere s, s1, s2; +float phi; + +uint * nObj; +float * prof; +Rayon * ray; +float3 * A, *u; +int t = 1; + + +void initObjet() +{ + srand(47); + node->s.r = 1.0f; + node[0].s.C = make_float3(0.0f,-1.5f,-0.0f); node[0].s.r = 0.5f; + node[1].s.C = make_float3(-1.0f,0.0f,-1.0f); node[1].s.r = 0.5f; + node[2].s.C = make_float3(1.0f,-0.f,-1.0f); node[2].s.r = 0.5f; + node[3].s.C = make_float3(0.0f,-0.f,-2.0f); node[3].s.r = 0.75f; + for( int i(4); i < numObj; i++ ) { + float r,v,b; + float tmp1(5.0f*((r=(float(rand()%255)/255.0f)))-2.5f); + float tmp2(5.0f*((v=(float(rand()%255)/255.0f)))-2.5f); + float tmp3(-5.0f*((b=(float(rand()%255)/255.0f)))); + float tmp4((rand()%100)/100.0f); + node[i].s.C = make_float3(tmp1,tmp2,tmp3); node[i].s.r = tmp4; + node[i].s.R = r; node[i].s.V = v; node[i].s.B = b; node[i].s.A = 1.0f; + node[i].fg = 0; node[i].fd = 0; + } + node[0].s.R = 0.0f; node[0].s.V = 1.0f; node[0].s.B = 1.0f; node[0].s.A = 1.0f; + node[1].s.R = 1.0f; node[1].s.V = 0.0f; node[1].s.B = 0.0f; node[1].s.A = 1.0f; + node[2].s.R = 0.0f; node[2].s.V = 0.0f; node[2].s.B = 1.0f; node[2].s.A = 1.0f; + node[3].s.R = 0.0f; node[3].s.V = 1.0f; node[3].s.B = 0.0f; node[3].s.A = 1.0f; + //createNode(&node[0], &node[1], &node[2], 1.0f); + node[0].fg = 1; node[0].fd = 2; + node[1].fg = 0; node[1].fd = 0; + node[2].fg = 0; node[2].fd = 0; + node[3].fg = 0; node[3].fd = 0; + + #ifdef DEBUG_RT_CUDA + h_debug_float4 = (float4*) calloc(DEBUG_NUM, sizeof(float4)); + h_debug_uint = (uint*) calloc(DEBUG_NUM, sizeof(uint)); + CUDA_SAFE_CALL( cudaMalloc( (void**)&d_debug_float4, DEBUG_NUM*sizeof(float4))); + CUDA_SAFE_CALL( cudaMalloc( (void**)&d_debug_uint, DEBUG_NUM*sizeof(uint))); + CUDA_SAFE_CALL( cudaMemcpy( d_debug_float4, h_debug_float4, DEBUG_NUM*sizeof(float4), cudaMemcpyHostToDevice) ); + CUDA_SAFE_CALL( cudaMemcpy( d_debug_uint, h_debug_uint, DEBUG_NUM*sizeof(uint), cudaMemcpyHostToDevice) ); + #endif + c_output = (uint*) calloc(width*height, sizeof(uint)); + CUDA_SAFE_CALL( cudaMalloc( (void**)&d_output, width*height*sizeof(uint))); + + CUDA_SAFE_CALL( cudaMalloc( (void**)&d_node, numObj*sizeof(Node) )); + CUDA_SAFE_CALL( cudaMemcpy( d_node, node, numObj*sizeof(Node), cudaMemcpyHostToDevice) ); + CUDA_SAFE_CALL( cudaMemcpyToSymbol(cnode, node, numObj*sizeof(Node)) ); + CUDA_SAFE_CALL( cudaMemcpyToSymbol(MView, (void*)&obs, 3*sizeof(float4)) ); + CUDA_SAFE_CALL( cudaMalloc( (void**)&d_temp, width * height*sizeof(uint))); + CUDA_SAFE_CALL( cudaMemset(d_temp, 0, width * height*sizeof(uint)) ); + + CUDA_SAFE_CALL( cudaMalloc( (void**)&nObj, width * height*sizeof(uint))); + CUDA_SAFE_CALL( cudaMalloc( (void**)&prof, width * height*sizeof(float))); + CUDA_SAFE_CALL( cudaMalloc( (void**)&ray, width * height*sizeof(Rayon))); + + CUDA_SAFE_CALL( cudaMalloc( (void**)&A, width * height*sizeof(float3))); + CUDA_SAFE_CALL( cudaMalloc( (void**)&u, width * height*sizeof(float3))); +} + +#define PRINT_PIXELS + +// Rendu de l'image avec CUDA +void render() +{ + // map PBO to get CUDA device pointer + //CUDA_SAFE_CALL(cudaGLMapBufferObject((void**)&d_output, pbo)); + //CUDA_SAFE_CALL( cudaMemcpy( d_output, c_output, width*height*sizeof(uint), cudaMemcpyHostToDevice) ); + // call CUDA kernel, writing results to PBO + CUT_SAFE_CALL(cutStartTimer(timer)); + #ifdef DEBUG_RT_CUDA + render<<>>(d_debug_float4, d_debug_uint, d_output, d_node, width, height, anim, obs.getDistance()); + #else + render<<>>(d_output, d_node, width, height, anim, obs.getDistance()); + #endif + CUDA_SAFE_CALL( cudaThreadSynchronize() ); + CUT_SAFE_CALL(cutStopTimer(timer)); + + #ifdef DEBUG_RT_CUDA + CUDA_SAFE_CALL( cudaMemcpy( h_debug_float4, d_debug_float4, DEBUG_NUM*sizeof(float4), cudaMemcpyDeviceToHost) ); + CUDA_SAFE_CALL( cudaMemcpy( h_debug_uint, d_debug_uint, DEBUG_NUM*sizeof(uint), cudaMemcpyDeviceToHost) ); + + printf("debug_float4\n"); + for (int i=0; i< DEBUG_NUM; i++) { + printf("%e %e %e %e\n", h_debug_float4[i].x, h_debug_float4[i].y, h_debug_float4[i].z, h_debug_float4[i].w); + } + printf("debug_uint\n"); + for (int i=0; i< DEBUG_NUM; i++) { + printf("0x%x\n", h_debug_uint[i]); + } + #endif + + CUDA_SAFE_CALL( cudaMemcpy( c_output, d_output, width*height*sizeof(uint), cudaMemcpyDeviceToHost) ); + unsigned long long int checksum = 0; + for (int y=(height-1); y >= 0; y--){ + if (g_verbose) printf("\n"); + for (int x=0; x< width; x++) { + if (g_verbose) printf("%010u ", (unsigned) c_output[x+y*width]); + checksum += c_output[x+y*width]; + } + } + printf("\n"); + printf("checksum=%llx\n", checksum); + CUT_CHECK_ERROR("Erreur kernel"); + + //CUDA_SAFE_CALL(cudaGLUnmapBufferObject(pbo)); // + +} + +// Affichage du resultat avec OpenGL +void display() +{ + + // Affichage du resultat + //glClear(GL_COLOR_BUFFER_BIT); + + //CUT_SAFE_CALL(cutStartTimer(timer)); + render(); + //CUT_SAFE_CALL(cutStopTimer(timer)); + printf("Kernel Time: %f \n", cutGetTimerValue(timer)); + /*fpsCount++; + if (fpsCount == fpsLimit) { + char fps[256]; + float ifps = 1.f / (cutGetAverageTimerValue(timer) / 1000.f); + sprintf(fps, "Cuda Ray Tracing: %.1f fps", ifps); + glutSetWindowTitle(fps); + fpsCount = 0; + fpsLimit = (int)max(ifps, 1.f); + CUT_SAFE_CALL(cutResetTimer(timer)); + }*/ + + if( anim >= 1.0f ) pas = -0.015f; + else if( anim <= -1.0f ) pas = 0.015f; + anim += pas; + + // Dessin de l'image de PBO + /*glDisable(GL_DEPTH_TEST); + glRasterPos2i(0, 0); + glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo); + glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, 0); + glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0); + + glutSwapBuffers(); + glutReportErrors();*/ + t--; + if (!t) { + return; + } + +} + +/*void idle() +{ + glutPostRedisplay(); +} + +void keyboard(unsigned char , int , int ) +{ + //glutPostRedisplay(); +}*/ + +int ox, oy; +int buttonState = 0; + +/*void mouse(int , int , int , int ) +{ + if (state == GLUT_DOWN) + buttonState |= 1<= 1) + break; + } + if (dev == deviceCount) { + fprintf(stderr, "There is no device supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + else + CUDA_SAFE_CALL(cudaSetDevice(dev)); + int i, commandline_error; + commandline_error = 0; + g_verbose = 0; + if (argc >= 3) { + width = atoi(argv[1]); + height = atoi(argv[2]); + for (i=3; i < argc;i++) { + if (argv[i][0] == '-') { + switch (argv[i][1]) { + case 'v': g_verbose = 1; + break; + default: commandline_error=1; + } + } + else commandline_error=1; + } + } else commandline_error=1; + + if (commandline_error || !width || !height) { + printf("Usage: ./rayTracing [-v]\n"); + printf("where WIDTH and HEIGHT are the screen dimensions and -v is used to display an abstract representation of the output.\n"); + return 1; + } + CUT_SAFE_CALL(cutCreateTimer(&timer)); + CUT_SAFE_CALL(cutResetTimer(timer)); + + initialize_bmp(width,height,32); + + // initialise les functions callback de GLUT + /*glutInit(&argc, argv); + glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE); + glutInitWindowSize(width, height); + glutCreateWindow("CUDA Ray Tracing"); + glutDisplayFunc(display); + glutKeyboardFunc(keyboard); + glutMouseFunc(mouse); + glutMotionFunc(motion); + glutReshapeFunc(reshape); + glutIdleFunc(idle); + + glewInit(); + if (!glewIsSupported("GL_VERSION_2_0 GL_ARB_pixel_buffer_object")) { + fprintf(stderr, "Les extensions minimales d'OpenGL sont absentes."); + exit(-1); + } + initPixelBuffer(); + initObjet(); + + atexit(cleanup); + + glutMainLoop();*/ + initObjet(); + initPixelBuffer(); + display(); + create_bmp(c_output); + CUT_SAFE_CALL(cutDeleteTimer(timer)); + return 0; +} -- cgit v1.3