LNK2005 errors

Hi devs.

I’ve just started experimenting with parallel computing. If I can get any help, I should get it here :)

I’m parallelizing a RayTracing algorithm using CUDA, and I am facing these errors (after solving hundreds of others):

LINK TO SCREENSHOT

I am using VS2010 with NSight, CUDA toolkit 4.2, GeForce 9500GS, driver v 306.32.

any ideas on how to get this running? I can also post the kernel and host code, if it helps.

Tnx!

So I tried moving everything into the .cu file, including host code and everything.

It compiles and runs successfuly, but the output is just a black screen, instead of the bouncing spheres. Otherwise I think it’s working, given that the times are shown (i’m measuring execution time).

Any ideas on how to solve this one and get the spheres to show?

Thank you

#define _USE_MATH_DEFINES
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include "C:\Users\MC.Choice\Documents

#define _USE_MATH_DEFINES
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include “C:\Users\MC.Choice\Documents\ŠOLA FAX\3. Letnik\Vzpredni in porazdeljeni sistemi in algoritmi\glew-1.10.0\glew-1.10.0\include\GL\glew.h”
#include <GL\glut.h>
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
#include <vector_types.h>
#define NSPHERES 4
#define VIEWPLANE 400
#define WINDOW VIEWPLANE*2
#define FOCALDIST 1000
#define RADIUS 200

typedef struct {
double r, g, b;
}pixel;

pixel *d_pixels,h_pixels;
#ifndef DEFS_H
#define DEFS_H
#define _BLINNPHONG
//#define _LAMBERT
//#define _NOSHADING
//typedef enum {false=0, true=1} BOOLEAN;
typedef enum {down = 0, up = 1} DIRECTION;
#define CRED 0
#define CGREEN 1
#define CBLUE 2
/
--------------- VECTORS -------------------- /
typedef struct {
double x;
double y;
double z;
}VECTOR3D;
/
----------------- VIEWPORT ----------------- /
typedef struct {
int xvmin;
int yvmin;
int xvmax;
int yvmax;
}VIEWPORT;
/
------------------- PIXEL ------------------ /
typedef struct {
int i;
int j;
}PIXEL;
/
---------------- SPHERE -------------------- */
typedef struct {
double lambda_in;
double lambda_out;
VECTOR3D normal;
VECTOR3D point;
bool valid;
}SPHERE_INTERSECTION;

typedef struct {
VECTOR3D center;
double radius;
double kd_rgb[3];
double ks_rgb[3];
double ka_rgb[3];
double shininess;
bool mirror;
}SPHERE;
/* ------------------- RAY --------------------- /
typedef struct {
VECTOR3D origin;
VECTOR3D direction;
}RAY;
/
--------------- VECTOR BASIS ---------------- */
typedef struct {
VECTOR3D u;
VECTOR3D v;
VECTOR3D n;
}VEC_BASIS;

#endif

device VEC_BASIS camera_frame;
device VECTOR3D view_point, light;
device SPHERE sphere[NSPHERES];
device VIEWPORT viewport;
device double focal_distance, color, light_intensity, ambi_light_intensity;
device DIRECTION direction[NSPHERES];
//vector.cpp
device void vec_sub(VECTOR3D *v1, VECTOR3D *v2, VECTOR3D *v3) {
v1->x = v2->x - v3->x;
v1->y = v2->y - v3->y;
v1->z = v2->z - v3->z;
}

device void vec_add(VECTOR3D *v1, VECTOR3D *v2, VECTOR3D *v3) {
v1->x = v2->x + v3->x;
v1->y = v2->y + v3->y;
v1->z = v2->z + v3->z;
}

device void vec_scale(double scale, VECTOR3D *v1, VECTOR3D *v2) {
v1->x = scale * v2->x;
v1->y = scale * v2->y;
v1->z = scale * v2->z;
}

device double dotproduct(VECTOR3D *v1, VECTOR3D *v2) {
return v1->x * v2->x + v1->y * v2->y + v1->z * v2->z;
}

device void normalize_vector(VECTOR3D *v) {
double magnitude;
magnitude = sqrt(dotproduct(v, v));
v->x = v->x / magnitude;
v->y = v->y / magnitude;
v->z = v->z / magnitude;
}
//vector.cpp
//sphere.cpp
device bool sphere_intersection (RAY *ray, SPHERE sphere, SPHERE_INTERSECTION intersection) {
double discriminant;
double A, B, C;
double lambda1, lambda2;
VECTOR3D temp;

A = dotproduct(&ray->direction, &ray->direction);

vec_sub(&temp, &ray->origin, &sphere->center);
B = 2 * dotproduct(&temp, &ray->direction);

vec_sub(&temp, &ray->origin, &sphere->center);
C = dotproduct(&temp, &temp) - (sphere->radius * sphere->radius);

discriminant = B*B - 4*A*C;

if (discriminant >= 0) {
	lambda1 = (-B + sqrt(discriminant)) / (2*A);
	lambda2 = (-B - sqrt(discriminant)) / (2*A);
	// is the object visible from the eye (lambda1,2>0)
	if (lambda1>=0 && lambda2>=0) {
		if (lambda1 == lambda2) {
			intersection->lambda_in = intersection->lambda_out = lambda1;
		} else if (lambda1 < lambda2) {
			intersection->lambda_in  = lambda1;
			intersection->lambda_out = lambda2;
		} else {
			intersection->lambda_in  = lambda2;
			intersection->lambda_out = lambda1;
		}
		intersection->valid = true;
		return true;
	} else {
		intersection->valid = false;
		return false;
	}
} else {
	intersection->valid = false;
	return false;
}

}

device void intersection_normal(SPHERE sphere, SPHERE_INTERSECTION intersection, RAY* ray) {
double lambda, scale;
VECTOR3D v1, v2, point, normal;

lambda = intersection->lambda_in;

vec_scale(lambda, &v1, &ray->direction);
vec_add(&point, &v1, &ray->origin);

intersection->point.x = point.x;
intersection->point.y = point.y;
intersection->point.z = point.z;

vec_sub(&v2, &point, &sphere->center);

scale = 1.0 / sphere->radius;
vec_scale(scale, &normal, &v2);

normalize_vector(&normal);

intersection->normal.x = normal.x;
intersection->normal.y = normal.y;
intersection->normal.z = normal.z;

}

device void copy_intersection_struct(SPHERE_INTERSECTION* is1, SPHERE_INTERSECTION* is2) {
is1->lambda_in = is2->lambda_in;
is1->lambda_out = is2->lambda_out;

is1->normal.x = is2->normal.x;
is1->normal.y = is2->normal.y;
is1->normal.z = is2->normal.z;

is1->point.x = is2->point.x;
is1->point.y = is2->point.y;
is1->point.z = is2->point.z;

is1->valid = is2->valid;

}
//ray.cpp
device void compute_ray(RAY* ray, VECTOR3D* view_point, VIEWPORT* viewport, PIXEL* pixel, VEC_BASIS* camera_frame, double distance) {
float u, v;
VECTOR3D v1, v2, v3, v4, dir;
// 1. calculate u and v coordinates of the pixels on the image plane:
u = (double)(viewport->xvmin) + (double)(pixel->i) + 0.5 ;
v = (double)(viewport->yvmin) + (double)(pixel->j) + 0.5 ;
// 2. calculate ray direction
vec_scale(-distance, &v1, &camera_frame->n);
vec_scale(u, &v2, &camera_frame->u);
vec_scale(v, &v3, &camera_frame->v);

ray->origin.x = view_point->x;
ray->origin.y = view_point->y;
ray->origin.z = view_point->z;

vec_add(&v4, &v1, &v2);
vec_add(&dir, &v4, &v3);
normalize_vector(&dir);

ray->direction.x = dir.x;
ray->direction.y = dir.y;
ray->direction.z = dir.z;

}

device void compute_reflected_ray(RAY* reflected_ray, RAY* incidence_ray, SPHERE_INTERSECTION* intersection) {
double dp1;
VECTOR3D scaled_normal, reflected_direction;
// calculate dot-product between surface normal and the direction of the incidence ray:
dp1 = dotproduct(&intersection->normal, &incidence_ray->direction);
// scale surface normal by 2dp1:
dp1 = 2
dp1;
vec_scale(dp1, &scaled_normal, &intersection->normal);

vec_sub(&reflected_direction, &incidence_ray->direction, &scaled_normal);

reflected_ray->origin.x = intersection->point.x;
reflected_ray->origin.y = intersection->point.y;
reflected_ray->origin.z = intersection->point.z;

reflected_ray->direction.x = reflected_direction.x;
reflected_ray->direction.y = reflected_direction.y;
reflected_ray->direction.z = reflected_direction.z;

}

device void compute_shadow_ray(RAY* ray, SPHERE_INTERSECTION* intersection, VECTOR3D* light) {
VECTOR3D dir;
// ray origin is in the intersection point
ray->origin.x = intersection->point.x;
ray->origin.y = intersection->point.y;
ray->origin.z = intersection->point.z;
// ray direction is from the intersection point towards the light:
vec_sub(&dir, light, &intersection->point);
normalize_vector(&dir);

ray->direction.x = dir.x;
ray->direction.y = dir.y;
ray->direction.z = dir.z;

}
//ray.cpp
//shading.cpp
device double lambertian_shading(SPHERE_INTERSECTION intersection, VECTOR3D light, double kd, double ka, double intensity, double amb_intensity) {
double color;
VECTOR3D l;
// compute vector l
vec_sub(&l, light, &intersection->point);
normalize_vector(&l);

// compute the intensity:
color = kd * intensity * dotproduct(&l, &intersection->normal) + ka * amb_intensity;
if (color > 0.0) {
	return color;
}
else {
	return 0.0;
}

}

device double blinnphong_shading(SPHERE_INTERSECTION intersection, VECTOR3D light, VECTOR3D* viewpoint, double kd, double ks,
double ka, double p, double intensity, double amb_intensity) {

double color_diffuse = 0.0;
double color_specular = 0.0;

VECTOR3D l;
VECTOR3D h;
VECTOR3D v;

// compute vector v :
vec_sub(&v, viewpoint, &intersection->point);
normalize_vector(&v);
// compute vector l :
vec_sub(&l, light, &intersection->point);
normalize_vector(&l);
// compute vector h:
vec_add(&h, &v, &l);
normalize_vector(&h);

// compute the diffuse intensity:
color_diffuse = kd * intensity * dotproduct(&l, &intersection->normal) ;
if (color_diffuse < 0.0) color_diffuse = 0.0;
// compute the specular intensity:
color_specular = ks * intensity * pow (dotproduct(&h, &intersection->normal), p);
if (color_specular < 0.0) color_specular = 0.0;

return (color_diffuse + color_specular + (ka * amb_intensity));

}

device double shadow(double ka, double amb_intensity) {
return (ka * amb_intensity);
}

device void set_rgb_array(double* rgb_array, double cred, double cgreen, double cblue) {
rgb_array[CRED] = cred;
rgb_array[CGREEN] = cgreen;
rgb_array[CBLUE] = cblue;
}
//shadow.cpp

void Timer(int obsolete) {
glutPostRedisplay();
glutTimerFunc(10, Timer, 0);
}

global void init(void) {
direction[0] = up;
direction[1] = down;
direction[2] = up;
// set scene:
viewport.xvmin = -VIEWPLANE;
viewport.yvmin = -VIEWPLANE;
viewport.xvmax = VIEWPLANE;
viewport.yvmax = VIEWPLANE;
//
camera_frame.u.x = 1.0;
camera_frame.u.y = 0.0;
camera_frame.u.z = 0.0;
//
camera_frame.v.x = 0.0;
camera_frame.v.y = 1.0;
camera_frame.v.z = 0.0;
//
camera_frame.n.x = 0.0;
camera_frame.n.y = 0.0;
camera_frame.n.z = 1.0;
//
view_point.x = (viewport.xvmax - viewport.xvmin) / 2.0 ;
view_point.y = (viewport.yvmax - viewport.yvmin) / 2.0 ;
view_point.z = 0.0;
//
light.x = view_point.x - 1300;
light.y = view_point.y + 1300 ;
light.z = view_point.z - 300;
//
ambi_light_intensity = 1.0;
light_intensity = 1.0;
//
focal_distance = FOCALDIST;
//
sphere[0].radius = RADIUS/1.5;
sphere[0].center.x = view_point.x - (RADIUS+30);
sphere[0].center.y = view_point.y ;
sphere[0].center.z = view_point.z - focal_distance - (2RADIUS+20);
// the first sphere is blue:
set_rgb_array(sphere[0].kd_rgb, 0.0, 0.0, 0.8);
set_rgb_array(sphere[0].ks_rgb, 1.0, 1.0, 1.0);
set_rgb_array(sphere[0].ka_rgb, 0.0, 0.0, 0.2);
sphere[0].shininess = 100.0;
sphere[0].mirror = false;
//
sphere[1].radius = RADIUS/1.2;
sphere[1].center.x = view_point.x + 0;
sphere[1].center.y = view_point.y + 50;
sphere[1].center.z = view_point.z - focal_distance - (3
RADIUS+20);
// the second sphere is green:
set_rgb_array(sphere[1].kd_rgb, 0.0, 0.5, 0.0);
set_rgb_array(sphere[1].ks_rgb, 1.0, 1.0, 1.0);
set_rgb_array(sphere[1].ka_rgb, 0.0, 0.2, 0.0);
sphere[1].shininess = 10.0;
sphere[1].mirror = false;
//
sphere[2].radius = RADIUS;
sphere[2].center.x = view_point.x + (2RADIUS+30);
sphere[2].center.y = view_point.y + 100;
sphere[2].center.z = view_point.z - focal_distance - (4
RADIUS+20);
// the third sphere is red:
set_rgb_array(sphere[2].kd_rgb, 1.0, 0.0, 0.0);
set_rgb_array(sphere[2].ks_rgb, 1.0, 1.0, 1.0);
set_rgb_array(sphere[2].ka_rgb, 0.2, 0.0, 0.0);
sphere[2].shininess = 100.0;
sphere[2].mirror = false;
//
sphere[3].radius = 100RADIUS;
sphere[3].center.x = view_point.x ;
sphere[3].center.y = view_point.y - 100
RADIUS-130;
sphere[3].center.z = view_point.z - focal_distance - (4*RADIUS+20);
// the third sphere is red:
set_rgb_array(sphere[3].kd_rgb, 0.5, 0.5, 0.5);
set_rgb_array(sphere[3].ks_rgb, 1.0, 1.0, 1.0);
set_rgb_array(sphere[3].ka_rgb, 0.5, 0.5, 0.5);
sphere[3].shininess = 100.0;
sphere[3].mirror = true;
}

global void animate(void) {
// bounce the first ball:
if (direction[0] == up) {
if (sphere[0].center.y < viewport.yvmax+300) sphere[0].center.y += 10;
else direction[0] = down;
} else if (direction[0] == down) {
if (sphere[0].center.y > viewport.yvmin+900) sphere[0].center.y -= 10;
else direction[0] = up;
}
// bounce the second ball:
if (direction[1] == up) {
if (sphere[1].center.y < viewport.yvmax+300) sphere[1].center.y += 20;
else direction[1] = down;
} else if (direction[1] == down) {
if (sphere[1].center.y > viewport.yvmin+900) sphere[1].center.y -= 20;
else direction[1] = up;
}
// bounce the third ball:
if (direction[2] == up) {
if (sphere[2].center.y < viewport.yvmax+300) sphere[2].center.y += 55;
else direction[2] = down;
} else if (direction[2] == down) {
if (sphere[2].center.y > viewport.yvmin+900) sphere[2].center.y -= 55;
else direction[2] = up;
}
}

global void calc(pixel *pixels) {
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= (viewport.xvmax - viewport.xvmin) || j > (viewport.yvmax - viewport.yvmin)) {
return;
} else {
int intersection_object = -1, k; // none
double theta, current_lambda = 0x7fefffffffffffff; // maximum positive double
RAY ray, shadow_ray;
PIXEL pixel;
SPHERE_INTERSECTION intersection, current_intersection, shadow_ray_intersection;
bool bShadow = false;
pixel.i = i;
pixel.j = j;
compute_ray(&ray, &view_point, &viewport, &pixel, &camera_frame, focal_distance);
for (k = 0; k < NSPHERES; k++) {
if (sphere_intersection(&ray, &sphere[k], &intersection)) {
intersection_normal(&sphere[k], &intersection, &ray);
if (intersection.lambda_in < current_lambda) {
current_lambda = intersection.lambda_in;
intersection_object = k;
copy_intersection_struct(&current_intersection, &intersection);
}
}
}
if (intersection_object > -1) {
compute_shadow_ray(&shadow_ray, &intersection, &light);
theta = dotproduct(&(shadow_ray.direction), &(intersection.normal));
for (k = 0; k < NSPHERES; k++) {
if (k != intersection_object) {
if (sphere_intersection(&shadow_ray, &sphere[k], &shadow_ray_intersection) && (theta > 0.0))
bShadow = true;
}
}
if(bShadow){
pixels[i * WINDOW + j].r = shadow(sphere[intersection_object].ka_rgb[CRED], ambi_light_intensity);
pixels[i * WINDOW + j].g = shadow(sphere[intersection_object].ka_rgb[CGREEN], ambi_light_intensity);
pixels[i * WINDOW + j].b = shadow(sphere[intersection_object].ka_rgb[CBLUE], ambi_light_intensity);
}else{
pixels[i * WINDOW + j].r = blinnphong_shading(&current_intersection, &light, &view_point,
sphere[intersection_object].kd_rgb[CRED], sphere[intersection_object].ks_rgb[CRED],
sphere[intersection_object].ka_rgb[CRED], sphere[intersection_object].shininess,
light_intensity, ambi_light_intensity);
pixels[i * WINDOW + j].g = blinnphong_shading(&current_intersection, &light, &view_point,
sphere[intersection_object].kd_rgb[CGREEN], sphere[intersection_object].ks_rgb[CGREEN],
sphere[intersection_object].ka_rgb[CGREEN], sphere[intersection_object].shininess,
light_intensity, ambi_light_intensity);
pixels[i * WINDOW + j].b = blinnphong_shading(&current_intersection, &light, &view_point,
sphere[intersection_object].kd_rgb[CBLUE], sphere[intersection_object].ks_rgb[CBLUE],
sphere[intersection_object].ka_rgb[CBLUE], sphere[intersection_object].shininess,
light_intensity, ambi_light_intensity);
}
} else {
pixels[i * WINDOW + j].r = 0;
pixels[i * WINDOW + j].g = 0;
pixels[i * WINDOW + j].b = 0;
}
}
}

void disp(void) {
cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);

//premik kugl
animate<<<1,1>>>();
cudaThreadSynchronize();
//clear all pixels:
glClear(GL_COLOR_BUFFER_BIT);
// RAY TRACING:
dim3 block(16,16, 1);
dim3 grid((WINDOW + (block.x - 1)) / block.x, (WINDOW + (block.y - 1)) / block.y, 1);
cudaEventRecord(start, 0);
calc<<<grid, block>>>(d_pixels);
cudaThreadSynchronize();
cudaEventRecord(stop, 0);
//Prenos rezultatov iz GPU v RAM
int velikost =  sizeof(pixel) * WINDOW * WINDOW;
cudaMemcpy((void*) h_pixels, (void*) d_pixels,velikost, cudaMemcpyDeviceToHost);
//Izris točk
int i, j;
for(i = 0; i < WINDOW; i++) {
	for(j = 0; j < WINDOW; j++) {
		glColor3f(h_pixels[i * WINDOW + j].r, h_pixels[i * WINDOW + j].g, h_pixels[i * WINDOW + j].b);
		glBegin(GL_POINTS);
		glVertex2i(i, j);
		glEnd();
	}
}
glFlush();
glutSwapBuffers();
cudaEventElapsedTime(&time, start, stop);
printf("%f\n", time);

}

int main(int argc, char** argv) {
//table for pixels value
h_pixels = (pixel*) malloc(sizeof(pixel) * WINDOW * WINDOW);
cudaMalloc((void**) &d_pixels, sizeof(pixel) * WINDOW * WINDOW);
// init glut:
glutInit (&argc, argv);
// specify the display mode to be RGB and single buffering:
glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);
// specify the initial window position:
glutInitWindowPosition(100, 100);
// specify the initial window size:
glutInitWindowSize(WINDOW, WINDOW);
// create the window and set title:
glutCreateWindow(“Basic Ray Tracer by Pa3cio, UL FRI”);
// init opengl:
init<<<1,1>>>();
glClearColor(0.0, 0.0, 0.0, 0.0);
glMatrixMode(GL_PROJECTION);
gluOrtho2D(0.0, WINDOW, 0.0, WINDOW);
// register callback function to display graphics:
glutDisplayFunc(disp);
// call Timer():
Timer(0);
// enter tha main loop and process events:
glutMainLoop();
//Free tables
free(h_pixels);
cudaFree(d_pixels);
return 0;
}

Scaron;OLA FAX. Letnik\Vzpredni in porazdeljeni sistemi in algoritmi\glew-1.10.0\glew-1.10.0\include\GL\glew.h"
#include <GL\glut.h>
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
#include <vector_types.h>
#define NSPHERES 4
#define VIEWPLANE 400
#define WINDOW VIEWPLANE*2
#define FOCALDIST 1000
#define RADIUS 200

typedef struct {
	double r, g, b;
}pixel;

pixel *d_pixels,*h_pixels;
#ifndef DEFS_H
#define DEFS_H
#define _BLINNPHONG
//#define _LAMBERT
//#define _NOSHADING
//typedef enum {false=0, true=1} BOOLEAN;
typedef enum {down = 0, up = 1} DIRECTION;
#define CRED 0
#define CGREEN 1
#define CBLUE 2
/* --------------- VECTORS -------------------- */
typedef struct {
	double x;
	double y;
	double z;
}VECTOR3D;
/* ----------------- VIEWPORT ----------------- */
typedef struct {
	int xvmin;
	int yvmin;
	int xvmax;
	int yvmax;
}VIEWPORT;
/* ------------------- PIXEL ------------------ */
typedef struct {
	int i;
	int j;
}PIXEL;
/* ---------------- SPHERE -------------------- */
typedef struct {
	double	lambda_in;
	double	lambda_out;
	VECTOR3D	normal;
	VECTOR3D point;
	bool	valid;
}SPHERE_INTERSECTION;

typedef struct {
	VECTOR3D center;
	double radius;
	double kd_rgb[3];
	double ks_rgb[3];
	double ka_rgb[3];
	double shininess;
	bool mirror;
}SPHERE;
/* ------------------- RAY --------------------- */
typedef struct {
	VECTOR3D origin;
	VECTOR3D direction;
}RAY;
/* --------------- VECTOR BASIS ---------------- */
typedef struct {
	VECTOR3D u;
	VECTOR3D v;
	VECTOR3D n;
}VEC_BASIS;

#endif

__device__ VEC_BASIS camera_frame;
__device__ VECTOR3D view_point, light;
__device__ SPHERE sphere[NSPHERES];
__device__ VIEWPORT viewport;
__device__ double focal_distance, color, light_intensity, ambi_light_intensity;
__device__ DIRECTION direction[NSPHERES];
//vector.cpp
__device__ void vec_sub(VECTOR3D *v1, VECTOR3D *v2, VECTOR3D *v3) {
	v1->x = v2->x - v3->x;
	v1->y = v2->y - v3->y;
	v1->z = v2->z - v3->z;
}

__device__ void vec_add(VECTOR3D *v1, VECTOR3D *v2, VECTOR3D *v3) {
	v1->x = v2->x + v3->x;
	v1->y = v2->y + v3->y;
	v1->z = v2->z + v3->z;
}

__device__ void vec_scale(double scale, VECTOR3D *v1, VECTOR3D *v2) {
	v1->x = scale * v2->x;
	v1->y = scale * v2->y;
	v1->z = scale * v2->z;
}

__device__ double dotproduct(VECTOR3D *v1, VECTOR3D *v2) {
	return v1->x * v2->x + v1->y * v2->y + v1->z * v2->z;
}

__device__ void normalize_vector(VECTOR3D *v) {
	double magnitude;
	magnitude = sqrt(dotproduct(v, v));
	v->x = v->x / magnitude;
	v->y = v->y / magnitude;
	v->z = v->z / magnitude;
}
//vector.cpp
//sphere.cpp
__device__ bool sphere_intersection (RAY *ray, SPHERE *sphere, SPHERE_INTERSECTION* intersection) {
	double discriminant;
	double A, B, C;
	double lambda1, lambda2;
	VECTOR3D temp;

	A = dotproduct(&ray->direction, &ray->direction);

	vec_sub(&temp, &ray->origin, &sphere->center);
	B = 2 * dotproduct(&temp, &ray->direction);

	vec_sub(&temp, &ray->origin, &sphere->center);
	C = dotproduct(&temp, &temp) - (sphere->radius * sphere->radius);

	discriminant = B*B - 4*A*C;

	if (discriminant >= 0) {
		lambda1 = (-B + sqrt(discriminant)) / (2*A);
		lambda2 = (-B - sqrt(discriminant)) / (2*A);
		// is the object visible from the eye (lambda1,2>0)
		if (lambda1>=0 && lambda2>=0) {
			if (lambda1 == lambda2) {
				intersection->lambda_in = intersection->lambda_out = lambda1;
			} else if (lambda1 < lambda2) {
				intersection->lambda_in  = lambda1;
				intersection->lambda_out = lambda2;
			} else {
				intersection->lambda_in  = lambda2;
				intersection->lambda_out = lambda1;
			}
			intersection->valid = true;
			return true;
		} else {
			intersection->valid = false;
			return false;
		}
	} else {
		intersection->valid = false;
		return false;
	}
}

__device__ void intersection_normal(SPHERE *sphere, SPHERE_INTERSECTION* intersection, RAY* ray) {
	double lambda, scale;
	VECTOR3D v1, v2, point, normal;

	lambda = intersection->lambda_in;

	vec_scale(lambda, &v1, &ray->direction);
	vec_add(&point, &v1, &ray->origin);

	intersection->point.x = point.x;
	intersection->point.y = point.y;
	intersection->point.z = point.z;

	vec_sub(&v2, &point, &sphere->center);

	scale = 1.0 / sphere->radius;
	vec_scale(scale, &normal, &v2);

	normalize_vector(&normal);

	intersection->normal.x = normal.x;
	intersection->normal.y = normal.y;
	intersection->normal.z = normal.z;

}

__device__ void copy_intersection_struct(SPHERE_INTERSECTION* is1, SPHERE_INTERSECTION* is2) {
	is1->lambda_in = is2->lambda_in;
	is1->lambda_out = is2->lambda_out;

	is1->normal.x = is2->normal.x;
	is1->normal.y = is2->normal.y;
	is1->normal.z = is2->normal.z;

	is1->point.x = is2->point.x;
	is1->point.y = is2->point.y;
	is1->point.z = is2->point.z;

	is1->valid = is2->valid;
}
//ray.cpp
__device__ void compute_ray(RAY* ray, VECTOR3D* view_point, VIEWPORT* viewport, PIXEL* pixel, VEC_BASIS* camera_frame, double distance) {
	float u, v;
	VECTOR3D v1, v2, v3, v4, dir;
	// 1. calculate u and v coordinates of the pixels on the image plane:
	u = (double)(viewport->xvmin) + (double)(pixel->i) + 0.5 ;
	v = (double)(viewport->yvmin) + (double)(pixel->j) + 0.5 ;
	// 2. calculate ray direction
	vec_scale(-distance, &v1, &camera_frame->n);
	vec_scale(u, &v2, &camera_frame->u);
	vec_scale(v, &v3, &camera_frame->v);

	ray->origin.x = view_point->x;
	ray->origin.y = view_point->y;
	ray->origin.z = view_point->z;

	vec_add(&v4, &v1, &v2);
	vec_add(&dir, &v4, &v3);
	normalize_vector(&dir);

	ray->direction.x = dir.x;
	ray->direction.y = dir.y;
	ray->direction.z = dir.z;
}

__device__ void compute_reflected_ray(RAY* reflected_ray, RAY* incidence_ray, SPHERE_INTERSECTION* intersection) {
	double dp1;
	VECTOR3D scaled_normal, reflected_direction;
	// calculate dot-product between surface normal and the direction of the incidence ray:
	dp1 = dotproduct(&intersection->normal, &incidence_ray->direction);
	// scale surface normal by 2*dp1:
	dp1 = 2*dp1;
	vec_scale(dp1, &scaled_normal, &intersection->normal);

	vec_sub(&reflected_direction, &incidence_ray->direction, &scaled_normal);

	reflected_ray->origin.x = intersection->point.x;
	reflected_ray->origin.y = intersection->point.y;
	reflected_ray->origin.z = intersection->point.z;

	reflected_ray->direction.x = reflected_direction.x;
	reflected_ray->direction.y = reflected_direction.y;
	reflected_ray->direction.z = reflected_direction.z;
}

__device__ void compute_shadow_ray(RAY* ray, SPHERE_INTERSECTION* intersection, VECTOR3D* light) {
	VECTOR3D dir;
	// ray origin is in the intersection point
	ray->origin.x = intersection->point.x;
	ray->origin.y = intersection->point.y;
	ray->origin.z = intersection->point.z;
	// ray direction is from the intersection point towards the light:
	vec_sub(&dir, light, &intersection->point);
	normalize_vector(&dir);

	ray->direction.x = dir.x;
	ray->direction.y = dir.y;
	ray->direction.z = dir.z;
}
//ray.cpp
//shading.cpp
__device__ double lambertian_shading(SPHERE_INTERSECTION *intersection, VECTOR3D* light, double kd, double ka, double intensity, double amb_intensity) {
	double color;
	VECTOR3D l;
	// compute vector l
	vec_sub(&l, light, &intersection->point);
	normalize_vector(&l);

	// compute the intensity:
	color = kd * intensity * dotproduct(&l, &intersection->normal) + ka * amb_intensity;
	if (color > 0.0) {
		return color;
	}
	else {
		return 0.0;
	}
}

__device__ double blinnphong_shading(SPHERE_INTERSECTION *intersection, VECTOR3D* light, VECTOR3D* viewpoint, double kd, double ks,
		double ka, double p, double intensity, double amb_intensity) {

	double color_diffuse = 0.0;
	double color_specular = 0.0;

	VECTOR3D l;
	VECTOR3D h;
	VECTOR3D v;

	// compute vector v :
	vec_sub(&v, viewpoint, &intersection->point);
	normalize_vector(&v);
	// compute vector l :
	vec_sub(&l, light, &intersection->point);
	normalize_vector(&l);
	// compute vector h:
	vec_add(&h, &v, &l);
	normalize_vector(&h);

	// compute the diffuse intensity:
	color_diffuse = kd * intensity * dotproduct(&l, &intersection->normal) ;
	if (color_diffuse < 0.0) color_diffuse = 0.0;
	// compute the specular intensity:
	color_specular = ks * intensity * pow (dotproduct(&h, &intersection->normal), p);
	if (color_specular < 0.0) color_specular = 0.0;

	return (color_diffuse + color_specular + (ka * amb_intensity));
}

__device__ double shadow(double ka, double amb_intensity) {
	return (ka * amb_intensity);
}

__device__ void set_rgb_array(double* rgb_array, double cred, double cgreen, double cblue) {
	rgb_array[CRED] = cred;
	rgb_array[CGREEN] = cgreen;
	rgb_array[CBLUE] = cblue;
}
//shadow.cpp

void Timer(int obsolete) {
	glutPostRedisplay();
	glutTimerFunc(10, Timer, 0);
}

__global__ void init(void) {
	direction[0] = up;
	direction[1] = down;
	direction[2] = up;
	// set scene:
	viewport.xvmin = -VIEWPLANE;
	viewport.yvmin = -VIEWPLANE;
	viewport.xvmax = VIEWPLANE;
	viewport.yvmax = VIEWPLANE;
	//
	camera_frame.u.x = 1.0;
	camera_frame.u.y = 0.0;
	camera_frame.u.z = 0.0;
	//
	camera_frame.v.x = 0.0;
	camera_frame.v.y = 1.0;
	camera_frame.v.z = 0.0;
	//
	camera_frame.n.x = 0.0;
	camera_frame.n.y = 0.0;
	camera_frame.n.z = 1.0;
	//
	view_point.x = (viewport.xvmax - viewport.xvmin) / 2.0 ;
	view_point.y = (viewport.yvmax - viewport.yvmin) / 2.0 ;
	view_point.z = 0.0;
	//
	light.x = view_point.x - 1300;
	light.y = view_point.y + 1300 ;
	light.z = view_point.z - 300;
	//
	ambi_light_intensity = 1.0;
	light_intensity = 1.0;
	//
	focal_distance = FOCALDIST;
	//
	sphere[0].radius = RADIUS/1.5;
	sphere[0].center.x  = view_point.x - (RADIUS+30);
	sphere[0].center.y  = view_point.y ;
	sphere[0].center.z  = view_point.z - focal_distance - (2*RADIUS+20);
	// the first sphere is blue:
	set_rgb_array(sphere[0].kd_rgb, 0.0, 0.0, 0.8);
	set_rgb_array(sphere[0].ks_rgb, 1.0, 1.0, 1.0);
	set_rgb_array(sphere[0].ka_rgb, 0.0, 0.0, 0.2);
	sphere[0].shininess = 100.0;
	sphere[0].mirror = false;
	//
	sphere[1].radius = RADIUS/1.2;
	sphere[1].center.x  = view_point.x + 0;
	sphere[1].center.y  = view_point.y + 50;
	sphere[1].center.z  = view_point.z - focal_distance - (3*RADIUS+20);
	// the second sphere is green:
	set_rgb_array(sphere[1].kd_rgb, 0.0, 0.5, 0.0);
	set_rgb_array(sphere[1].ks_rgb, 1.0, 1.0, 1.0);
	set_rgb_array(sphere[1].ka_rgb, 0.0, 0.2, 0.0);
	sphere[1].shininess = 10.0;
	sphere[1].mirror = false;
	//
	sphere[2].radius = RADIUS;
	sphere[2].center.x  = view_point.x + (2*RADIUS+30);
	sphere[2].center.y  = view_point.y + 100;
	sphere[2].center.z  = view_point.z - focal_distance - (4*RADIUS+20);
	// the third sphere is red:
	set_rgb_array(sphere[2].kd_rgb, 1.0, 0.0, 0.0);
	set_rgb_array(sphere[2].ks_rgb, 1.0, 1.0, 1.0);
	set_rgb_array(sphere[2].ka_rgb, 0.2, 0.0, 0.0);
	sphere[2].shininess = 100.0;
	sphere[2].mirror = false;
	//
	sphere[3].radius = 100*RADIUS;
	sphere[3].center.x  = view_point.x ;
	sphere[3].center.y  = view_point.y - 100*RADIUS-130;
	sphere[3].center.z  = view_point.z - focal_distance - (4*RADIUS+20);
	// the third sphere is red:
	set_rgb_array(sphere[3].kd_rgb, 0.5, 0.5, 0.5);
	set_rgb_array(sphere[3].ks_rgb, 1.0, 1.0, 1.0);
	set_rgb_array(sphere[3].ka_rgb, 0.5, 0.5, 0.5);
	sphere[3].shininess = 100.0;
	sphere[3].mirror = true;
}

__global__ void animate(void) {
	// bounce the first ball:
	if (direction[0] == up) {
		if (sphere[0].center.y < viewport.yvmax+300) sphere[0].center.y += 10;
		else direction[0] = down;
	} else if (direction[0] == down) {
		if (sphere[0].center.y > viewport.yvmin+900) sphere[0].center.y -= 10;
		else direction[0] = up;
	}
	// bounce the second ball:
	if (direction[1] == up) {
		if (sphere[1].center.y < viewport.yvmax+300) sphere[1].center.y += 20;
		else direction[1] = down;
	} else if (direction[1] == down) {
		if (sphere[1].center.y > viewport.yvmin+900) sphere[1].center.y -= 20;
		else direction[1] = up;
	}
	// bounce the third ball:
	if (direction[2] == up) {
		if (sphere[2].center.y < viewport.yvmax+300) sphere[2].center.y += 55;
		else direction[2] = down;
	} else if (direction[2] == down) {
		if (sphere[2].center.y > viewport.yvmin+900) sphere[2].center.y -= 55;
		else direction[2] = up;
	}
}

__global__ void calc(pixel *pixels) {
	unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int j = blockIdx.y * blockDim.y + threadIdx.y;
	if (i >= (viewport.xvmax - viewport.xvmin) || j > (viewport.yvmax - viewport.yvmin)) {
		return;
	} else {
		int intersection_object = -1, k; // none
		double theta, current_lambda = 0x7fefffffffffffff; // maximum positive double
		RAY ray, shadow_ray;
		PIXEL pixel;
		SPHERE_INTERSECTION intersection, current_intersection, shadow_ray_intersection;
		bool bShadow = false;
		pixel.i = i;
		pixel.j = j;
		compute_ray(&ray, &view_point, &viewport, &pixel, &camera_frame, focal_distance);
		for (k = 0; k < NSPHERES; k++) {
			if (sphere_intersection(&ray, &sphere[k], &intersection)) {
				intersection_normal(&sphere[k], &intersection, &ray);
				if (intersection.lambda_in < current_lambda) {
					current_lambda = intersection.lambda_in;
					intersection_object = k;
					copy_intersection_struct(&current_intersection, &intersection);
				}
			}
		}
		if (intersection_object > -1) {
			compute_shadow_ray(&shadow_ray, &intersection, &light);
			theta = dotproduct(&(shadow_ray.direction), &(intersection.normal));
			for (k = 0; k < NSPHERES; k++) {
				if (k != intersection_object)	{
					if (sphere_intersection(&shadow_ray, &sphere[k], &shadow_ray_intersection) && (theta > 0.0))
						bShadow = true;
				}
			}
			if(bShadow){
				pixels[i * WINDOW + j].r = shadow(sphere[intersection_object].ka_rgb[CRED], ambi_light_intensity);
				pixels[i * WINDOW + j].g = shadow(sphere[intersection_object].ka_rgb[CGREEN], ambi_light_intensity);
				pixels[i * WINDOW + j].b = shadow(sphere[intersection_object].ka_rgb[CBLUE], ambi_light_intensity);
			}else{
				pixels[i * WINDOW + j].r = blinnphong_shading(&current_intersection, &light, &view_point,
					sphere[intersection_object].kd_rgb[CRED], sphere[intersection_object].ks_rgb[CRED],
					sphere[intersection_object].ka_rgb[CRED], sphere[intersection_object].shininess,
					light_intensity, ambi_light_intensity);
				pixels[i * WINDOW + j].g = blinnphong_shading(&current_intersection, &light, &view_point,
					sphere[intersection_object].kd_rgb[CGREEN], sphere[intersection_object].ks_rgb[CGREEN],
					sphere[intersection_object].ka_rgb[CGREEN], sphere[intersection_object].shininess,
					light_intensity, ambi_light_intensity);
				pixels[i * WINDOW + j].b = blinnphong_shading(&current_intersection, &light, &view_point,
					sphere[intersection_object].kd_rgb[CBLUE], sphere[intersection_object].ks_rgb[CBLUE],
					sphere[intersection_object].ka_rgb[CBLUE], sphere[intersection_object].shininess,
					light_intensity, ambi_light_intensity);
			}
		} else {
			pixels[i * WINDOW + j].r = 0;
			pixels[i * WINDOW + j].g = 0;
			pixels[i * WINDOW + j].b = 0;
		}
	}
}

void disp(void) {
	cudaEvent_t start, stop;
	float time;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	
	//premik kugl
	animate<<<1,1>>>();
	cudaThreadSynchronize();
	//clear all pixels:
	glClear(GL_COLOR_BUFFER_BIT);
	// RAY TRACING:
	dim3 block(16,16, 1);
	dim3 grid((WINDOW + (block.x - 1)) / block.x, (WINDOW + (block.y - 1)) / block.y, 1);
	cudaEventRecord(start, 0);
	calc<<<grid, block>>>(d_pixels);
	cudaThreadSynchronize();
	cudaEventRecord(stop, 0);
	//Prenos rezultatov iz GPU v RAM
	int velikost =  sizeof(pixel) * WINDOW * WINDOW;
	cudaMemcpy((void*) h_pixels, (void*) d_pixels,velikost, cudaMemcpyDeviceToHost);
	//Izris točk
	int i, j;
	for(i = 0; i < WINDOW; i++) {
		for(j = 0; j < WINDOW; j++) {
			glColor3f(h_pixels[i * WINDOW + j].r, h_pixels[i * WINDOW + j].g, h_pixels[i * WINDOW + j].b);
			glBegin(GL_POINTS);
			glVertex2i(i, j);
			glEnd();
		}
	}
	glFlush();
	glutSwapBuffers();
	cudaEventElapsedTime(&time, start, stop);
	printf("%f\n", time);
}

int main(int argc, char** argv) {
	//table for pixels value
	h_pixels = (pixel*) malloc(sizeof(pixel) * WINDOW * WINDOW);
	cudaMalloc((void**) &d_pixels, sizeof(pixel) * WINDOW * WINDOW);
	// init glut:
	glutInit (&argc, argv);
	// specify the display mode to be RGB and single buffering:
	glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);
	// specify the initial window position:
	glutInitWindowPosition(100, 100);
	// specify the initial window size:
	glutInitWindowSize(WINDOW, WINDOW);
	// create the window and set title:
	glutCreateWindow("Basic Ray Tracer by Pa3cio, UL FRI");
	// init opengl:
	init<<<1,1>>>();
	glClearColor(0.0, 0.0, 0.0, 0.0);
	glMatrixMode(GL_PROJECTION);
	gluOrtho2D(0.0, WINDOW, 0.0, WINDOW);
	// register callback function to display graphics:
	glutDisplayFunc(disp);
	// call Timer():
	Timer(0);
	// enter tha main loop and process events:
	glutMainLoop();
	//Free tables
	free(h_pixels);
	cudaFree(d_pixels);
	return 0;
}