# Accelerated Ray Tracing in One Weekend in CUDA

**URL:** <https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615>\
**Category:** Technical Blog\
**Created:** [November 5, 2018, 2:12pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615 "2018-11-05T14:12:08Z")\
**Posts on this page:** 20\
**Page:** 1

<div class="post-metadata">

**Author:** ![jwitsoe](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/jwitsoe/32/16591_2.png) [@jwitsoe](https://forums.developer.nvidia.com/u/jwitsoe)\
**Post date:** [November 5, 2018, 2:12pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/1 "2018-11-05T14:12:08Z")

</div>

Originally published at: [https://developer.nvidia.com/blog/accelerated-ray-tracing-cuda/](https://developer.nvidia.com/blog/accelerated-ray-tracing-cuda/)  
   
Recent announcements of NVIDIA’s new Turing GPUs, RTX technology, and Microsoft’s DirectX Ray Tracing have spurred a renewed interest in ray tracing. Using these technologies vastly simplifies the ability to write applications using ray tracing. But what if you’re curious about how ray tracing actually works? One way to learn is to code your own ray tracing…

---

<div class="post-metadata">

**Author:** ![anon88161431](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon88161431](https://forums.developer.nvidia.com/u/anon88161431)\
**Post date:** [November 6, 2018, 2:29am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/2 "2018-11-06T02:29:08Z")

</div>

Great Post just wanted to point out a typo you made at chapter 12.

"Using a new GeForce RTX 1080, the image renders in just 4.7 seconds for a 19x speedup!"

Was kinda funny :)

---

<div class="post-metadata">

**Author:** ![anon941035](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon941035](https://forums.developer.nvidia.com/u/anon941035)\
**Post date:** [November 6, 2018, 5:20pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/3 "2018-11-06T17:20:22Z")

</div>

Good catch! We fixed it in the text, thanks.

---

<div class="post-metadata">

**Author:** ![anon39217652](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon39217652](https://forums.developer.nvidia.com/u/anon39217652)\
**Post date:** [November 10, 2018, 12:09pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/4 "2018-11-10T12:09:44Z")

</div>

What is fb? You keep referencing fb in your code but it is undefined when I try to run the code. I am just trying to get the code to run and show graphics.

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [November 13, 2018, 2:01am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/5 "2018-11-13T02:01:57Z")

</div>

Thank you for your question. It was an oversight to not show how I allocate the frame buffer (aka "fb") that holds the pixel RGB values. I’ll work with the editor to update the post. This is shown in the Chapter 1 branch of the reference code here: [https://github.com/rogerall...](https://github.com/rogerallen/raytracinginoneweekendincuda/blob/ch01_output_cuda/main.cu#L39-L41 "https://github.com/rogerallen/raytracinginoneweekendincuda/blob/ch01\_output\_cuda/main.cu#L39-L41")

---

<div class="post-metadata">

**Author:** ![anon39217652](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon39217652](https://forums.developer.nvidia.com/u/anon39217652)\
**Post date:** [November 14, 2018, 8:36am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/6 "2018-11-14T08:36:25Z")

</div>

Thank you for your contribution. It is greatly appreciated by the CUDA community!

---

<div class="post-metadata">

**Author:** ![anon39217652](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon39217652](https://forums.developer.nvidia.com/u/anon39217652)\
**Post date:** [November 17, 2018, 6:13pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/7 "2018-11-17T18:13:59Z")

</div>

I was able to run your code successfully! On my CPU it takes 30 minutes to generate the image, but with GPU assistance using CUDA I'm down to approx 3 minutes. That's a really really big deal! I have a nvidia GeForce 1050 ti graphics card.

One last thing in regards to the code. On the later chapters, the CUDA code uses M\_PI constant in the camera.h code. Please add the following snippet before the camera.h include so there are no build-time errors:

#ifndef M\_PI  
#define M\_PI 3.14159265358979323846  
#endif

Again, thanks for your contribution. I will tinker with the code so I can learn ray tracing better. Have a great day!

-CUDA Education

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [November 19, 2018, 4:43pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/8 "2018-11-19T16:43:22Z")

</div>

I'm happy to hear you were successful in translating the code & your suggestion is a good one. I have found that I need that #define if I try out the code on Windows. Thanks!

---

<div class="post-metadata">

**Author:** ![anon45886813](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon45886813](https://forums.developer.nvidia.com/u/anon45886813)\
**Post date:** [January 27, 2019, 9:53pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/9 "2019-01-27T21:53:21Z")

</div>

Great post!  
Somehow on my GTX 1050 and CUDA 9.2, my machine freezes completely at the end of the rendering. After a lot of trials I figured it happens when accessing fb on cpu while writing the image to disk. Copying the buffer to a host array resolved the issue for me.  
Any ideas what could be causing this ?

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [February 5, 2019, 5:53am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/10 "2019-02-05T05:53:33Z")

</div>

The only thing I can think of is that you somehow removed the cudaDeviceSynchronize() after the render call & before the first read from fb. That could cause an issue like this.

If that wasn't it, suggest you try to find the simplest version of the code that doesn't work (try going back to Chapter 1 to see if that passes) and we could attempt to debug via an Issue on the github repository.

---

<div class="post-metadata">

**Author:** ![anon45886813](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon45886813](https://forums.developer.nvidia.com/u/anon45886813)\
**Post date:** [February 6, 2019, 2:31pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/11 "2019-02-06T14:31:03Z")

</div>

Let my try to simplify the code and still reproduce the issue, will open a github issue when I do. Thanks again!

---

<div class="post-metadata">

**Author:** ![anon91724277](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon91724277](https://forums.developer.nvidia.com/u/anon91724277)\
**Post date:** [June 6, 2019, 11:29am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/12 "2019-06-06T11:29:57Z")

</div>

Thanks very much for this !! I was hoping someone could give some insight into some issues I'm having in running this using VS 2017 on Windows10. I basically started a blank CUDA 9.1 Runtime project in VS 2017, then copied and pasted the header files and [main.cu](http://main.cu "main.cu") from the Chapter 12 code. I also selected the VS 2015 v140 platform, and under CUDA C/C++/Device in the Configuration Properties changed Code Generation to "compute\_60,sm\_60" to presumably match my GTX 1080ti. I then ran the debugger, and aside from the Intellisense warnings for the kernel launcher ("\<\<\< \>\>\>") and not recognizing "blockidx", it ran fine. And I sent the output to a .ppm file and dragged that into Gimp and the rendered image looked fine.

My problem is that it seems that it never utilized the GTX-1080ti, since the render (10 samples) took about 46 seconds, and Task Manager showed no significant GPU activity with either the 3D or CUDA (or any other) engines. Instead it seemed to bounce between using 2 of my CPU cores (8 - core Ryzen).

My other concern is that it never opened a window showing the realtime GPU render as it progressed, only printing the RGB values in the command window.

I'm guessing I missed a step in configuring the GPU in Visual Studio. Anyone have any insights? Thanks much.

Oh, and if anyone has any additional insight into adding a simple UI to this like ImGUI so I can manually tweak settings during the render I'd be thrilled....

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [June 14, 2019, 10:43pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/13 "2019-06-14T22:43:27Z")

</div>

I am not sure what could be happening with your VS 2017 project. If you see CUDA calls in the code it will use the GPU--there's no way around that. I've done something similar and gotten sub-10-second speeds with a GTX 1070. Maybe debug string output is super-slow? Make sure you create a Release executable would be one suggestion.

Note that there is no GUI for this project. It only outputs to the console.

---

<div class="post-metadata">

**Author:** ![anon91724277](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon91724277](https://forums.developer.nvidia.com/u/anon91724277)\
**Post date:** [June 14, 2019, 11:04pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/14 "2019-06-14T23:04:08Z")

</div>

BINGO !!!! 4.5 seconds !!! It was the Release vs. Debug as you mentioned. DOH !!! Thanks much.

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [June 15, 2019, 1:19am UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/15 "2019-06-15T01:19:18Z")

</div>

Fantastic!

---

<div class="post-metadata">

**Author:** ![anon4486737](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon4486737](https://forums.developer.nvidia.com/u/anon4486737)\
**Post date:** [July 2, 2019, 6:53pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/16 "2019-07-02T18:53:11Z")

</div>

Dear Roger Allen,

Thank you very much  
for your excellent implementation or ray tracing/casting in CUDA. I  
was really surprised that you manged to use virtual functions from  
and the abstract class hitable to create other classes. You manged to  
derive hitable\_list and sphere classes from the hitable abstract  
class. All manuals on CUDA published by NVIDIA claim that it is  
impossible to use virtual functions and abstract classes in order to  
create other classes for usage in global kernel functions. Following  
you steps I derived from hitable class other classes, like Plane,   
Instance (for moving and rotation), Compound (for drawing complex  
objects). All of them can be nested in each other. I even managed to  
implement the class Grid for acceleration of computations called  
regular grids described in the book Ray Tracing from the Ground up by  
Keven Suffern. You employed the global kernel called Create World in  
order to fill in GPU memory with objects. But when the number of  
object was greater than 4000 spheres packed in the acceleration class  
Grid I received the mistake

CUDA error = 700 at  
[main5\_add\_obj\_grid\_init.cu](http://main5_add_obj_grid_init.cu "main5\_add\_obj\_grid\_init.cu"):256 'cudaDeviceSynchronize()'   
(program exited with  
code: 99)

Then I decided to  
use unified memory cudaMallocManaged((void\*\*)&d\_world,  
sizeof(hitable \*)); and employed host function to instead of the  
kernel function create\_world and immediately received the error  
message   
number of objects=34  
CUDA error = 700 at  
[main5\_add\_obj\_grid\_init.cu](http://main5_add_obj_grid_init.cu "main5\_add\_obj\_grid\_init.cu"): 270 'cudaDeviceSynchronize()'

I have Quadro P4000  
card with 8 gigabytes of memory. My host device has 32 gigabytes. I  
use Ubuntu operating system and CUDA 10.1  
Potentially, one can  
use host functions to access and modify unified memory and then use  
global kernel functions to process that unified memory . If I am  
right my comp has in total 40 gigabytes of unified memory. I was  
able to ray trace even several millions of spheres on my modest  
laptop with 8 gigabytes without any CUDA card for this purpose. So ,  
what was wrong in my case?  
Best Regards,  
Valery Pavlov,  
Spain

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [July 12, 2019, 5:19pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/17 "2019-07-12T17:19:40Z")

</div>

Hi Valery,

I am going to have to look closer into my usage of virtual functions and the CUDA documentation and I will try to find some internal experts to help me understand. It was an oversight and not my intention to go against a CUDA rule. It looks like I may have stumbled into some undefined behavior. So, it may well be the case that my usage does not trigger errors, but your more elaborate usage does trigger errors as you seem to have found.

Now that I have read some posts like [this one](https://devtalk.nvidia.com/default/topic/1026116/cuda-programming-and-performance/copying-objects-to-device-with-virtual-functions/post/5218666/#5218666 "https://devtalk.nvidia.com/default/topic/1026116/cuda-programming-and-performance/copying-objects-to-device-with-virtual-functions/post/5218666/#5218666"), I could see how the second case of creating objects on the host and copying them to the device would lose virtual function information and cause crashes.

For debugging your first issue that creates the objects on the device, since you have made significant changes, I think the best plan would be for you to post to the [CUDA Programming forum](https://devtalk.nvidia.com/default/board/57/cuda-programming-and-performance "https://devtalk.nvidia.com/default/board/57/cuda-programming-and-performance") so that others can try to reproduce and investigate the error.

---

<div class="post-metadata">

**Author:** ![anon4486737](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon4486737](https://forums.developer.nvidia.com/u/anon4486737)\
**Post date:** [July 22, 2019, 2:51pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/18 "2019-07-22T14:51:15Z")

</div>

Hi Roger I created the following topic:

[https://devtalk.nvidia.com/...](https://devtalk.nvidia.com/default/topic/1057750/cuda-programming-and-performance/ray-tracing-over-weekend-with-cuda-using-unified-memory/ "https://devtalk.nvidia.com/default/topic/1057750/cuda-programming-and-performance/ray-tracing-over-weekend-with-cuda-using-unified-memory/")

in which I provide the small modification of your code provided by you for chapter 5:

Dear all ,

I have read a very  
interesting post Accelerated Ray Tracing in One Weekend in CUDA by By  
Roger Allen  
([https://devblogs.nvidia.com...](https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/) "https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/)").  
In this post virtual functions were employed to create ray-traced  
pictures of spheres. I modified the code provided in Chapter 5   
([https://github.com/rogerall...](https://github.com/rogerallen/raytracinginoneweekendincuda/tree/ch05_normals_cuda) "https://github.com/rogerallen/raytracinginoneweekendincuda/tree/ch05\_normals\_cuda)")

a little bit in  
order to create\_world on host instead of device. In other words I  
replaced \_\_device\_\_ functions with \_\_host\_\_ \_\_device\_\_ functions   
in headers ray.h; hitable.h; hitable\_list.h; sphere.h; and finally  
I added host function in [main.cu](http://main.cu "main.cu") file

void  
create\_world\_host(hitable \*\*d\_list, hitable \*\*d\_world) {  
\*(d\_list)= new sphere(vec3(0,0,-1), 0.5);

\*(d\_list+1)  
= new sphere(vec3(0,-100.5,-1), 100);

\*d\_world   
= new hitable\_list(d\_list,2);

}

and used unified  
memory

// make our  
world of hitables

//but using  
unified memory

hitable  
\*\*d\_list;

checkCudaErrors(cudaMallocManaged((void \*\*)&d\_list,  
2\*sizeof(hitable \*)));

hitable  
\*\*d\_world;

checkCudaErrors(cudaMallocManaged((void \*\*)&d\_world,  
sizeof(hitable \*)));)

I found it is  
impossible to fill in the unified memory using my host function   
create\_world\_host instead of similar \_\_global\_\_ void create\_world . I  
have Quadro P4000 card with Pascal architecture.  
Why does usage of  
the host function create\_world\_host instead of \_\_global\_\_ void  
create\_world lead to the error : endering  
a 1200x600 image in 8x8 blocks.  
CUDA error = 700 at  
[main.cu](http://main.cu "main.cu"):114 'cudaDeviceSynchronize()' ????

My version of ray.h  
is

#ifndef RAYH  
#define RAYH  
#include "vec3.h"  
class ray  
{

public:

\_\_host\_\_  
\_\_device\_\_ ray() {}

\_\_host\_\_  
\_\_device\_\_ ray(const vec3& a, const vec3& b) { A = a; B = b;  
}

\_\_host\_\_  
\_\_device\_\_ vec3 origin() const { return A; }

\_\_host\_\_  
\_\_device\_\_ vec3 direction() const { return B; }

\_\_host\_\_  
\_\_device\_\_ vec3 point\_at\_parameter(float t) const { return A + t\*B; }

vec3 A;

vec3 B;

};

#endif

My version of  
hitable.h is

#ifndef HITABLEH

#define HITABLEH

#include "ray.h"

struct hit\_record

{

float t;

vec3 p;

vec3 normal;

};

class hitable {

public:

\_\_host\_\_  
\_\_device\_\_ virtual bool hit(const ray& r, float t\_min, float  
t\_max, hit\_record& rec) const = 0;

};

#endif

My version of  
hitable\_list.h is:

#ifndef  
HITABLELISTH

#define HITABLELISTH

#include "hitable.h"

class hitable\_list:  
public hitable {

public:

\_\_host\_\_  
\_\_device\_\_ hitable\_list() {}

\_\_host\_\_  
\_\_device\_\_ hitable\_list(hitable \*\*l, int n) {list = l; list\_size = n;  
}

\_\_host\_\_  
\_\_device\_\_ virtual bool hit(const ray& r, float tmin, float tmax,  
hit\_record& rec) const;

hitable  
\*\*list;

int  
list\_size;

};

\_\_host\_\_ \_\_device\_\_  
bool hitable\_list::hit(const ray& r, float t\_min, float t\_max,  
hit\_record& rec) const {

hit\_record  
temp\_rec;

bool  
hit\_anything = false;

float  
closest\_so\_far = t\_max;

for (int i =  
0; i \< list\_size; i++) {

if  
(list[i]-\>hit(r, t\_min, closest\_so\_far, temp\_rec)) {

hit\_anything = true;

closest\_so\_far = temp\_rec.t;

rec  
= temp\_rec;

}

}

return  
hit\_anything;

}

#endif

My version of  
sphere.h is:

#ifndef SPHEREH

#define SPHEREH

#include "hitable.h"

class sphere: public  
hitable {

public:

\_\_host\_\_  
\_\_device\_\_ sphere() {}

\_\_host\_\_  
\_\_device\_\_ sphere(vec3 cen, float r) : center(cen), radius(r) {};

\_\_host\_\_  
\_\_device\_\_ virtual bool hit(const ray& r, float tmin, float tmax,  
hit\_record& rec) const;

vec3 center;

float  
radius;

};

\_\_host\_\_ \_\_device\_\_  
bool sphere::hit(const ray& r, float t\_min, float t\_max,  
hit\_record& rec) const {

vec3 oc =  
r.origin() - center;

float a =  
dot(r.direction(), r.direction());

float b =  
dot(oc, r.direction());

float c =  
dot(oc, oc) - radius\*radius;

float  
discriminant = b\*b - a\*c;

if (discriminant  
\> 0) {

float temp =  
(-b - sqrt(discriminant))/a;

if (temp \<  
t\_max && temp \> t\_min) {

rec.t =  
temp;

rec.p =  
r.point\_at\_parameter(rec.t);

rec.normal = (rec.p - center) / radius;

return  
true;

}

temp = (-b +  
sqrt(discriminant)) / a;

if (temp \<  
t\_max && temp \> t\_min) {

rec.t =  
temp;

rec.p =  
r.point\_at\_parameter(rec.t);

rec.normal = (rec.p - center) / radius;

return  
true;

}

}

return false;

}

#endif

and finally my  
version of [main.cu](http://main.cu "main.cu"):

#include \<iostream\>

#include \<time.h\>

#include \<float.h\>

#include "vec3.h"

#include "ray.h"

#include "sphere.h"

#include  
"hitable\_list.h"

// limited version  
of checkCudaErrors from helper\_cuda.h in CUDA examples

#define  
checkCudaErrors(val) check\_cuda( (val), #val, \_\_FILE\_\_, \_\_LINE\_\_ )

void  
check\_cuda(cudaError\_t result, char const \*const func, const char  
\*const file, int const line) {

if (result) {

std::cerr \<\<  
"CUDA error = " \<\< static\_cast\<unsigned int=""\>(result) \<\< " at " \<\<

file \<\<  
":" \<\< line \<\< " '" \<\< func \<\<  
"' \n";

// Make sure  
we call CUDA Device Reset before exiting

cudaDeviceReset();

exit(99);

}

}

\_\_device\_\_ vec3  
color(const ray& r, hitable \*\*world) {

hit\_record rec;

if  
((\*world)-\>hit(r, 0.0, FLT\_MAX, rec)) {

return  
0.5f\*vec3(rec.normal.x()+1.0f, rec.normal.y()+1.0f,  
rec.normal.z()+1.0f);

}

else {

vec3  
unit\_direction = unit\_vector(r.direction());

float t =  
0.5f\*(unit\_direction.y() + 1.0f);

return  
(1.0f-t)\*vec3(1.0, 1.0, 1.0) + t\*vec3(0.5, 0.7, 1.0);

}

}

\_\_global\_\_ void  
render(vec3 \*fb, int max\_x, int max\_y,

vec3 lower\_left\_corner, vec3 horizontal, vec3 vertical, vec3  
origin,

hitable \*\*world) {

int i =  
threadIdx.x + blockIdx.x \* blockDim.x;

int j =  
threadIdx.y + blockIdx.y \* blockDim.y;

if((i \>=  
max\_x) || (j \>= max\_y)) return;

int pixel\_index  
= j\*max\_x + i;

float u =  
float(i) / float(max\_x);

float v =  
float(j) / float(max\_y);

ray r(origin,  
lower\_left\_corner + u\*horizontal + v\*vertical);

fb[pixel\_index]  
= color(r, world);

}

\_\_global\_\_ void  
create\_world(hitable \*\*d\_list, hitable \*\*d\_world) {

if (threadIdx.x  
== 0 && blockIdx.x == 0) {

\*(d\_list)   
= new sphere(vec3(0,0,-1), 0.5);

\*(d\_list+1)  
= new sphere(vec3(0,-100.5,-1), 100);

\*d\_world   
= new hitable\_list(d\_list,2);

}

}

void  
create\_world\_host(hitable \*\*d\_list, hitable \*\*d\_world) {

\*(d\_list)   
= new sphere(vec3(0,0,-1), 0.5);

\*(d\_list+1)  
= new sphere(vec3(0,-100.5,-1), 100);

\*d\_world   
= new hitable\_list(d\_list,2);

}

\_\_global\_\_ void  
free\_world(hitable \*\*d\_list, hitable \*\*d\_world) {

delete  
\*(d\_list);

delete  
\*(d\_list+1);

delete \*d\_world;

}

int main() {

int nx = 1200;

int ny = 600;

int tx = 8;

int ty = 8;

std::cerr \<\<  
"Rendering a " \<\< nx \<\< "x" \<\<  
ny \<\< " image ";

std::cerr \<\<  
"in " \<\< tx \<\< "x" \<\< ty \<\<  
" blocks.\n";

int num\_pixels =  
nx\*ny;

size\_t fb\_size =  
num\_pixels\*sizeof(vec3);

// allocate FB

vec3 \*fb;

checkCudaErrors(cudaMallocManaged((void \*\*)&fb, fb\_size));

// make our  
world of hitables

//but using  
unified memory

hitable  
\*\*d\_list;

checkCudaErrors(cudaMallocManaged((void \*\*)&d\_list,  
2\*sizeof(hitable \*)));

hitable  
\*\*d\_world;

checkCudaErrors(cudaMallocManaged((void \*\*)&d\_world,  
sizeof(hitable \*)));

//\*\*\*Creation of  
world on device using unified memory works\*\*\*//

create\_world\<\<\<1,1\>\>\>(d\_list,d\_world);

//\*\*\*Creation of  
world on host using unified memory fails\*\*\*//

//  
create\_world\_host(d\_list,d\_world);

checkCudaErrors(cudaGetLastError());

checkCudaErrors(cudaDeviceSynchronize());

clock\_t start,  
stop;

start = clock();

// Render our  
buffer

dim3  
blocks(nx/tx+1,ny/ty+1);

dim3  
threads(tx,ty);

render\<\<\<blocks, threads=""\>\>\>(fb, nx, ny,

vec3(-2.0, -1.0, -1.0),

vec3(4.0, 0.0, 0.0),

vec3(0.0, 2.0, 0.0),

vec3(0.0, 0.0, 0.0),

d\_world);

checkCudaErrors(cudaGetLastError());

checkCudaErrors(cudaDeviceSynchronize());

stop = clock();

double  
timer\_seconds = ((double)(stop - start)) / CLOCKS\_PER\_SEC;

std::cerr \<\<  
"took " \<\< timer\_seconds \<\< " seconds.\n";

// Output FB as  
Image

std::cout \<\<  
"P3\n" \<\< nx \<\< " " \<\< ny \<\<  
"\n255\n";

for (int j =  
ny-1; j \>= 0; j--) {

for (int i =  
0; i \< nx; i++) {

size\_t  
pixel\_index = j\*nx + i;

int ir =  
int(255.99\*fb[pixel\_index].r());

int ig =  
int(255.99\*fb[pixel\_index].g());

int ib =  
int(255.99\*fb[pixel\_index].b());

std::cout \<\< ir \<\< " " \<\< ig \<\< "  
" \<\< ib \<\< "\n";

}

}

// clean up

checkCudaErrors(cudaDeviceSynchronize());

free\_world\<\<\<1,1\>\>\>(d\_list,d\_world);

checkCudaErrors(cudaGetLastError());

checkCudaErrors(cudaFree(d\_list));

checkCudaErrors(cudaFree(d\_world));

checkCudaErrors(cudaFree(fb));

// useful for  
cuda-memcheck --leak-check full

cudaDeviceReset();

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [July 22, 2019, 3:48pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/19 "2019-07-22T15:48:15Z")

</div>

I'm happy to see this is already getting replies. Suggest further followups go there for clarity.

---

<div class="post-metadata">

**Author:** ![anon76518154](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@anon76518154](https://forums.developer.nvidia.com/u/anon76518154)\
**Post date:** [July 23, 2019, 3:53pm UTC](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615/20 "2019-07-23T15:53:53Z")

</div>

To clarify, I did discuss a bit with some internal experts and found for certain that my code is not relying on any undefined behavior (UB).

I was worried about this statement:

> It is not allowed to pass as an argument to a \_\_global\_\_ function **an object** of a class with virtual functions.

from [https://docs.nvidia.com/cud...](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#virtual-functions "https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#virtual-functions")

Passing an object by value could trigger that UB, but since I am passing a **pointer** to an object that is both created & used on the device, the code works as expected. Internally, we are also discussing how to best update the programming guide for better clarity.

Further discussion should probably go on the devtalk forum post that Valery created in his response to the previous post.

[Next page](https://forums.developer.nvidia.com/t/accelerated-ray-tracing-in-one-weekend-in-cuda/148615.md?page=2)