<p>It is <strong>easy</strong> to write a parallel program, but <strong>difficult</strong> to make it fast, just as you've said: naive porting might just be bad. If you think your GPU renderer could be faster but actually it is not and if there is no mistake made by you, then it will be the problem of your device code design: each SM might not even compete with individual CPU core when the parallelism is poorly implemented.
Though ray (path) tracing is seemingly easy to run in parallel since computation is independent for each pixel, there are still several things that can make it slower:</p>
<ul>
<li>Warp divergence. Obviously, the device function <code>trace</code> will lead to a tremendous amount of warp divergence (since you have multiple branches, and for loops (even, nested)), causing your warps of threads to be serialized. Note that, CPU is a more powerful device for intense logic and branching and when your kernel has tons of possible branching, it can be slow. Try run <code>ncu --section WarpStateStats (your CUDA program)</code> and see the profiling results about warp efficiency. The logic in a thread should avoid branching as much as possible, otherwise it would violate the idea of SIMT. You said in the post that you got rid of the branches induced by different materials, but it seems to me that it's far from enough.</li>
<li>Why not use <code>float</code> instead of <code>double</code>? <code>float</code> is not that inaccurate, even in some extreme cases (for example, multiple scattering in volumetric rendering). The performance difference can be dramatic (greater than 2x, see <a href="https://www.reddit.com/r/cpp_questions/comments/1bjl3d5/why_cuda_c_uses_float/" rel="nofollow noreferrer">here</a>), and some of the instructions does not even support <code>double</code> (for example, <code>atomicAdd</code> supports <code>float</code> but not <code>double</code>, AFAIK). You may think that CPU uses <code>double</code>, too so it should be a fair match but no: GPU has far more ALUs than CPU, therefore the hardware design should be simpler, meaning that each ALU in GPU should be less competent than its CPU counterpart, let alone the modern CPU are almost all 64-bits.</li>
<li>In case you don't know. Simple scenes like this can be stored (partially) in shared memory to accelerate memory transaction: what's being reused a lot should be moved to faster memory. It is easily noticed that your code, in its current state, has much memory bound ops, aside from nested <code>for</code> loops and unpredictable branches: you are not doing any heavy math ops like evaluating normal distribution function in microfacet models, and just reading from your primitive positions and do intersection tests. Either use constant memory (for its independent constant cache, which might bring the cache efficiency higher) or shared memory.</li>
</ul>
<p>Actually, the necessary step is the <code>double</code> to <code>float</code> step. However, due to your scene settings (too big, and a bad design choice to use a huge sphere as flat surface, which leads to numerical instability when solving the equation in a large scene), I have to tweak your code a bit to make it work correctly (basically, scale the scenes of yours down by a factor of <code>0.00001</code>). You can see the results below, I am using a RTX TITAN GPU and a 112 core server scale CPU. With <code>double</code>, your renderer is indeed slower than CPU implementation, but with <code>float</code> type:</p>
<div class="s-table-container"><table class="s-table">
<thead>
<tr>
<th style="text-align: center;">float GPU (0.093s for ratio = 1.5) / (0.042s for ratio = 1)</th>
<th style="text-align: center;">float CPU  (1.502s for ratio = 1.5) / (0.687s for ratio = 1)</th>
</tr>
</thead>
<tbody>
<tr>
<td style="text-align: center;"><a href="https://i.sstatic.net/j06DC.jpg" rel="nofollow noreferrer"><img src="https://i.sstatic.net/j06DC.jpg" alt="enter image description here" /></a></td>
<td style="text-align: center;"><a href="https://i.sstatic.net/QoV4J.jpg" rel="nofollow noreferrer"><img src="https://i.sstatic.net/QoV4J.jpg" alt="enter image description here" /></a></td>
</tr>
</tbody>
</table></div>
<p>After this, I think you can experiment with some other acceleration tricks I've mentioned. The modified GPU verision can be found <a href="https://gist.github.com/Enigmatisms/801eca388493d6b6161b5ce917d19bec" rel="nofollow noreferrer">here</a>. Hope this helps.</p>

CC BY-SA 4.0
2024-04-13T01:52:53.937Z
Double precision atomics are available in CUDA since the Pascal architecture, see the [CUDA C++ Programming Guide](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd).
paleonix
CC BY-SA 4.0
2024-04-13T01:55:57.500Z
Thanks for the tip bro.
Enigmatisms
CC BY-SA 4.0
2024-04-12T08:47:05.383Z
2024-04-12T08:47:05.383Z
Enigmatisms
<p>Summary: I ported a simple path tracer <a href="https://www.kevinbeason.com/smallpt/" rel="nofollow noreferrer">smallpt</a> into CUDA but the performance doesn't boost.</p>
<ol>
<li>reproduce results:</li>
</ol>
<ul>
<li>build and run the <a href="https://github.com/w3ntao/smallpt-cpu" rel="nofollow noreferrer">CPU version</a>:</li>
</ul>
<pre><code>$ git clone --recurse-submodules https://github.com/w3ntao/smallpt-cpu.git

# build
<span class="math-container">$ cd smallpat-cpu
$</span> mkdir build; cd build
$ cmake ..; make -j

# render
$ ./smallpt-cpu
rendering (64 spp) took 10.600 seconds.
image saved to `smallpt_cpu_64.png`
</code></pre>
<ul>
<li>build and run the <a href="https://github.com/w3ntao/smallpt-megakernel" rel="nofollow noreferrer">CUDA version</a>:</li>
</ul>
<pre><code>$ git clone --recurse-submodules https://github.com/w3ntao/smallpt-megakernel.git

# build
<span class="math-container">$ cd smallpt-megakernel
$</span> mkdir build; cd build
$ cmake ..; make -j

# render
$ ./smallpt-megakernel
rendering (64 spp) took 11.214 seconds.
image saved to `smallpt_megakernel_64.png`
</code></pre>
<p>Both programs render visually correct results but for the same number of samples, it took around 10 sec for the CPU ray tracer to render and around 11 sec for the CUDA version.</p>
<p>My laptop setup: Intel Core i7-10750H (6 cores, 12 threads), NVIDIA GeForce GTX 1650</p>
<p><a href="https://i.sstatic.net/QMMTo.jpg" rel="nofollow noreferrer"><img src="https://i.sstatic.net/QMMTo.jpg" alt="enter image description here" /></a></p>
<ol start="2">
<li>core function <code>trace()</code> (in <code>main.cpp</code> and <code>main.cu</code>)</li>
</ol>
<p>This function is not complex, its pseudocode can be summarized as:</p>
<pre class="lang-cpp prettyprint-override"><code>Vec3 trace(Ray ray) {
    Vec3 radiance(0.0, 0.0, 0.0);
    Vec3 throughput(1.0, 1.0, 1.0);

    Ray ray &lt;- camera_ray;
    for (int depth = 0;; ++depth) {
        hit_point, sphere = intersect(ray, scene)
        // compute hit_point and sphere by intersecting the ray with the scene

        if (hit_nothing) {
            // stop the loop if hit nothing
            break;
        }

        update(radiance, sphere.emission)        // update radiance according to sphere's emission
        update(throughput, sphere.material_type) // update throughput according to sphere's material_type

        extension_ray = spawn_ray(ray, hit_point, sphere.material_type)
        // compute extension_ray from previous ray, hit_point and sphere.material_type

        ray &lt;- extension_ray
        // start next round bouncing with extension_ray
    }

    return radiance;
}
</code></pre>
<p>There are some mathematical and trigonometric details behind it but I don't think they are that much relevant in current discussion.</p>
<p>Other than different memory management, the code are virtually identical for CPU and CUDA implementation. For CPU, each thread would render a bunch of pixels, roughly 1/N of all pixels for a N-threads CPU. For GPU, each thread render 1 pixel, and block dimension is 8x8.</p>
<p>First I thought (different) branches in code could hurt kernel funcion performance, so I change 2 spheres (a mirror and a glass) to diffuse materials (<a href="https://github.com/w3ntao/smallpt-megakernel/blob/f81d7753ad38447bcac4330478cb996bc77ec3d6/src/main.cu#L332" rel="nofollow noreferrer">main.cu: line 332</a> and <a href="https://github.com/w3ntao/smallpt-megakernel/blob/f81d7753ad38447bcac4330478cb996bc77ec3d6/src/main.cu#L336" rel="nofollow noreferrer">main.cu: line 336</a>), rendering time improved to just 9.9 second (from 11 second).</p>
<p><a href="https://i.sstatic.net/xYITa.jpg" rel="nofollow noreferrer"><img src="https://i.sstatic.net/xYITa.jpg" alt="enter image description here" /></a>
I understand that such naive porting (from CPU functions) to CUDA is not optimized, but still I believe it with such sub-optimized solution, the CUDA ray tracer should still outperform the CPU version. So I'm asking what have I possibly done wrong to get this ray tracer rendering in such low performance?</p>

CC BY-SA 4.0
2024-04-12T06:57:48.037Z
1. Are you sure the code is actually running on the GPU (/CUDA)? Not sure how you did your port, but if both CPU and GPU are possible, it might be that the wrong path is chosen. 2. Why not port it to OptiX? It's made specifically for ray tracing. Even if your graphics card does not have ray tracing cores, this should be the better choice and I assume it needs less sync between CPU and GPU.
Tare
CC BY-SA 4.0
2024-04-12T07:11:01.783Z
@Tare 1. I'm sure the code is actually running on the GPU, it was running in a __global__ function. 2. I chose CUDA instead of Optix/OpenGL/Vulkan because I want to build everything from scratch to learn more about ray tracing.
Rahn
<raytracing><cuda>
Why is my ray tracer not accelerated by CUDA?
Rahn
# Why is my ray tracer not accelerated by CUDA?

(Asked by: username_0 on 2024-04-12)

Summary: I ported a simple path tracer [smallpt](https://www.kevinbeason.com/smallpt/) into CUDA but the performance doesn't boost.

1. reproduce results:

* build and run the [CPU version](https://github.com/w3ntao/smallpt-cpu):

```
$ git clone --recurse-submodules https://github.com/w3ntao/smallpt-cpu.git

# build
$ cd smallpat-cpu
$ mkdir build; cd build
$ cmake ..; make -j

# render
$ ./smallpt-cpu
rendering (64 spp) took 10.600 seconds.
image saved to `smallpt_cpu_64.png`

```

* build and run the [CUDA version](https://github.com/w3ntao/smallpt-megakernel):

```
$ git clone --recurse-submodules https://github.com/w3ntao/smallpt-megakernel.git

# build
$ cd smallpt-megakernel
$ mkdir build; cd build
$ cmake ..; make -j

# render
$ ./smallpt-megakernel
rendering (64 spp) took 11.214 seconds.
image saved to `smallpt_megakernel_64.png`

```

Both programs render visually correct results but for the same number of samples, it took around 10 sec for the CPU ray tracer to render and around 11 sec for the CUDA version.

My laptop setup: Intel Core i7-10750H (6 cores, 12 threads), NVIDIA GeForce GTX 1650

[![enter image description here](https://i.sstatic.net/QMMTo.jpg)](https://i.sstatic.net/QMMTo.jpg)

2. core function `trace()` (in `main.cpp` and `main.cu`)

This function is not complex, its pseudocode can be summarized as:

```cpp
Vec3 trace(Ray ray) {
    Vec3 radiance(0.0, 0.0, 0.0);
    Vec3 throughput(1.0, 1.0, 1.0);

    Ray ray <- camera_ray;
    for (int depth = 0;; ++depth) {
        hit_point, sphere = intersect(ray, scene)
        // compute hit_point and sphere by intersecting the ray with the scene

        if (hit_nothing) {
            // stop the loop if hit nothing
            break;
        }

        update(radiance, sphere.emission)        // update radiance according to sphere's emission
        update(throughput, sphere.material_type) // update throughput according to sphere's material_type

        extension_ray = spawn_ray(ray, hit_point, sphere.material_type)
        // compute extension_ray from previous ray, hit_point and sphere.material_type

        ray <- extension_ray
        // start next round bouncing with extension_ray
    }

    return radiance;
}

```

There are some mathematical and trigonometric details behind it but I don't think they are that much relevant in current discussion.

Other than different memory management, the code are virtually identical for CPU and CUDA implementation. For CPU, each thread would render a bunch of pixels, roughly 1/N of all pixels for a N-threads CPU. For GPU, each thread render 1 pixel, and block dimension is 8x8.

First I thought (different) branches in code could hurt kernel funcion performance, so I change 2 spheres (a mirror and a glass) to diffuse materials ([main.cu: line 332](https://github.com/w3ntao/smallpt-megakernel/blob/f81d7753ad38447bcac4330478cb996bc77ec3d6/src/main.cu#L332) and [main.cu: line 336](https://github.com/w3ntao/smallpt-megakernel/blob/f81d7753ad38447bcac4330478cb996bc77ec3d6/src/main.cu#L336)), rendering time improved to just 9.9 second (from 11 second).

[![enter image description here](https://i.sstatic.net/xYITa.jpg)](https://i.sstatic.net/xYITa.jpg)
I understand that such naive porting (from CPU functions) to CUDA is not optimized, but still I believe it with such sub-optimized solution, the CUDA ray tracer should still outperform the CPU version. So I'm asking what have I possibly done wrong to get this ray tracer rendering in such low performance?

**Tags:** raytracing, cuda

> **Comments:**
> **username_1**: 1. Are you sure the code is actually running on the GPU (/CUDA)? Not sure how you did your port, but if both CPU and GPU are possible, it might be that the wrong path is chosen. 2. Why not port it to OptiX? It's made specifically for ray tracing. Even if your graphics card does not have ray tracing cores, this should be the better choice and I assume it needs less sync between CPU and GPU.  
> **username_0**: @username_1 1. I'm sure the code is actually running on the GPU, it was running in a __global__ function. 2. I chose CUDA instead of Optix/OpenGL/Vulkan because I want to build everything from scratch to learn more about ray tracing.  

---

## 1 answers

### Answer by username_2 on 2024-04-12. Score: 3 | Accepted answer

It is **easy** to write a parallel program, but **difficult** to make it fast, just as you've said: naive porting might just be bad. If you think your GPU renderer could be faster but actually it is not and if there is no mistake made by you, then it will be the problem of your device code design: each SM might not even compete with individual CPU core when the parallelism is poorly implemented.
Though ray (path) tracing is seemingly easy to run in parallel since computation is independent for each pixel, there are still several things that can make it slower:

* Warp divergence. Obviously, the device function `trace` will lead to a tremendous amount of warp divergence (since you have multiple branches, and for loops (even, nested)), causing your warps of threads to be serialized. Note that, CPU is a more powerful device for intense logic and branching and when your kernel has tons of possible branching, it can be slow. Try run `ncu --section WarpStateStats (your CUDA program)` and see the profiling results about warp efficiency. The logic in a thread should avoid branching as much as possible, otherwise it would violate the idea of SIMT. You said in the post that you got rid of the branches induced by different materials, but it seems to me that it's far from enough.
* Why not use `float` instead of `double`? `float` is not that inaccurate, even in some extreme cases (for example, multiple scattering in volumetric rendering). The performance difference can be dramatic (greater than 2x, see [here](https://www.reddit.com/r/cpp_questions/comments/1bjl3d5/why_cuda_c_uses_float/)), and some of the instructions does not even support `double` (for example, `atomicAdd` supports `float` but not `double`, AFAIK). You may think that CPU uses `double`, too so it should be a fair match but no: GPU has far more ALUs than CPU, therefore the hardware design should be simpler, meaning that each ALU in GPU should be less competent than its CPU counterpart, let alone the modern CPU are almost all 64-bits.
* In case you don't know. Simple scenes like this can be stored (partially) in shared memory to accelerate memory transaction: what's being reused a lot should be moved to faster memory. It is easily noticed that your code, in its current state, has much memory bound ops, aside from nested `for` loops and unpredictable branches: you are not doing any heavy math ops like evaluating normal distribution function in microfacet models, and just reading from your primitive positions and do intersection tests. Either use constant memory (for its independent constant cache, which might bring the cache efficiency higher) or shared memory.

Actually, the necessary step is the `double` to `float` step. However, due to your scene settings (too big, and a bad design choice to use a huge sphere as flat surface, which leads to numerical instability when solving the equation in a large scene), I have to tweak your code a bit to make it work correctly (basically, scale the scenes of yours down by a factor of `0.00001`). You can see the results below, I am using a RTX TITAN GPU and a 112 core server scale CPU. With `double`, your renderer is indeed slower than CPU implementation, but with `float` type:

| float GPU (0.093s for ratio = 1.5) / (0.042s for ratio = 1) | float CPU (1.502s for ratio = 1.5) / (0.687s for ratio = 1) |
| --- | --- |
| [enter image description here](https://i.sstatic.net/j06DC.jpg) | [enter image description here](https://i.sstatic.net/QoV4J.jpg) |

After this, I think you can experiment with some other acceleration tricks I've mentioned. The modified GPU verision can be found [here](https://gist.github.com/username_2/801eca388493d6b6161b5ce917d19bec). Hope this helps.

> **Comments:**
> **username_3**: Double precision atomics are available in CUDA since the Pascal architecture, see the [CUDA C++ Programming Guide](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd).  
> **username_2**: Thanks for the tip bro.  

---

