GPUOpen

As the other posts have said, this isn't the right algorithm. Onesweep and its kin are cool but intimidating. The magic is easier to understand if one looks into the core algorithm: radix sort. A very readable explanation is here: https://gpuopen.com/download/publications/Introduction_to_GPU_Radix_So… It turns out that radix sort can be implemented in a way that is very straightforward to parallelize. It's a beautiful and elegant approach and worth knowing about!

Ok, good to know. To wit, I found the following: AMD: https://gpuopen.com/amd-gpu-architecture-programming-documentation/ Intel: https://www.intel.com/content/www/us/en/docs/graphics-for-linux/develo…

A 2% to 14% improvement in performance across the board is indeed good news - https://openbenchmarking.org/result/2112248-TJ-1080PRYZE71 ... Wish more developers would support and focus on opensource projects like AMD-GPU - https://gpuopen.com/ - instead of wasting their time trying to reverse-engineer a closed black box like the Apple M1.

It may be coming - there's a great slide deck on the RDNA changes and how to get the most out of the architecture: https://gpuopen.com/wp-content/uploads/2019/08/RDNA_Architecture_publi…

Tensorflow has a working AMD fork: https://github.com/ROCmSoftwarePlatform/tensorflow-upstream See following links for more discussion: https://www.reddit.com/r/MachineLearning/comments/85iwyj/d_tensorflow_… https://gpuopen.com/rocm-tensorflow-1-8-release/ https://stackoverflow.com/questions/37892784/using-keras-tensorflow-wi…

GPUOpen [1] is AMD's branded Open Source / Open protocol initiative. [1] https://gpuopen.com/

> better deep learning support They released the cuDNN equivalent for AMD devices a few weeks ago: http://gpuopen.com/developer-quick-start-miopen-1-0/ So AMD devices should now be capable of accelerating tensorflow, caffe, etc. (insert your favorite GPU deep learning framework here)

This adds more to the evidence that Vulkan / DX12 seems like a failed API design, when too many graphics engineers are reinventing the wheel by building their render-graph API on top of these (propietary triple-A game engines, Unreal Engine, and now Godot...) Instead of real-time graphics APIs providing all these low-level manual synchronization primitives and cumbersome PSOs, maybe they should just provide an official Render Graph API instead? Provide all of the operations and its dependencies in an acyclic graph up-front, and the driver handles synchronization automatically in the most performant manner tailored to the hardware. I guess there needed some trial-and-error in the gamedev world for about a decade to really nail down a nice to use but also performant graphics API design. Vulkan being originated from the Mantle API from AMD didn't help - since it was a low-level console API mainly accustomed to AMD's GPUs and really didn't seem like it would fit for a more "general-purpose" API spanning a huge range of hardware and can stand the test of time. And with Microsoft's DX12 hastely copying from AMD's initial design it also has all the same issues (The irony is that DX11 is still the best graphics API you can use in gamedev in Windows in terms of ergonomics and even performance - seeing many trying to dauntingly build a DX12 backend and end up performing worse than DX11...) Nowadays I'm obversing that the industry has known these issues for a while and are…

> What about generating foliage or varied ground textures? What about generating buildings? Or NPCs? This isn't how DLSS works at all. DLSS can't create new buildings or foliage, or NPCs, so hard to foresee problems of the kind Valve are concerned with. Same for varied ground textures- the entire point of the technology is to sharpen and upscale the original image, or in case of DLSS3 inject new matching frames. The only "risk" to DLSS is in Nvidia's own training data, but there is no "risk" of another company's existing IP leaking into final frame - again if there was, gamers wouldn't want to use it, as its destroying the original frame! If the resulting frame isn't a near perfect match for the original, DLSS has failed. Thankfully it does nearly perfectly match the original in use and alongside AMD's FSR 2.0 stuff has been one of the best advancements in gaming technology of recent years - effectively significant FPS boost "for free" on same hardware. > https://www.nvidia.com/en-us/geforce/technologies/dlss/ > https://gpuopen.com/fidelityfx-superresolution-2/

> Do you have a hardcore introduction with even more detail / perhaps even with examples of implementations? Hard to beat the actual technical references when you want something hardcore! The actual CUDA documentation is good: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html In particular, "Section 5: Performance Guidelines": https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#perfo… gives a lot of micro-architectural details, including those nasty "bank conflicts" that people keep talking about. Honestly, the original documentation says it best with just a few paragraphs on that particular manner: > To achieve high bandwidth, shared memory is divided into equally-sized memory modules, called banks, which can be accessed simultaneously. Any memory read or write request made of n addresses that fall in n distinct memory banks can therefore be serviced simultaneously, yielding an overall bandwidth that is n times as high as the bandwidth of a single module. > > However, if two addresses of a memory request fall in the same memory bank, there is a bank conflict and the access has to be serialized. The hardware splits a memory request with bank conflicts into as many separate conflict-free requests as necessary, decreasing throughput by a factor equal to the number of separate memory requests. If the number of separate memory requests is n, the initial memory request is said to cause n-way bank conflicts. See? Its really not…

I recently looked into GPGPU programming, and it seems to be a bit of a mess, which is surprising seeing how long it's been around and how much it is used. I think a large part of it is due to Nvidia using their strong position to hamper anything other than Cuda. OpenCL 2.0 was announced in 2013, Nvidia added OpenCL 2.0 "for evaluation purposes" in 2017. I don't think they have final support yet. OpenCL 2.1 uses the same intermediate language as Vulkan, SPIR-V. Nvidia does not seem to support this. When OpenCL 2.2 was announced in 2017 Khronos said they were working on converging OpenCL and Vulkan compute, but OpenCL will remain separate. I think this means that they will both use the same SPIR-V backend to run shader code, but I don't know. Khronos also has SYCL. The latest version, 1.2.1, was released a few days ago and uses OpenCL 1.2 as the backend. I'm guessing it doesn't use a newer version because of Nvidia's poor support for them, but AMD doesn't seem to support it well either. Then there's a multitude of other libraries/platforms, like POCL, HCC, Acoran, that I don't know anything about but found mentions of. AMD has HIP which can convert Cuda to run on AMD hardware, so maybe Cuda is the best option for AMD as well? I find it a bit of an embarrassment for the industry that it's this messy and complicated, and I don't see it getting better in the near future. It seems as if the proprietary Cuda is still the best option, and that's a big failure for everyone…

OpenGL still has lots of steam in it and 4.6 was recently released, so it isn't going away anytime soon. Even on mobile devices, Vulkan is hardly being embraced by Android developers, because it requires Android 7 or later, while being an optional API, which implies a tiny portion of their customer base. So very few devices support it, to the point Google doesn't think it is worthwhile to display on Android dashboards. https://developer.android.com/about/dashboards/ http://vulkan.gpuinfo.org/vulkansupport.php#android_devices AMD is also busy trying to cater CAD vendors to adopt Vulkan, with initiatives like V-EZ. https://gpuopen.com/v-ez-brings-easy-mode-vulkan/ However, if you want to be an engine developer, probably it would be better to have Vulkan knowledge than the APIs that everyone else already knows. Some guys at Khronos sessions (check their YouTube) stated they would favor candidates in interviews in such cases. For product development, I am a bit fan of middleware where each API is just a tick box away, though.

Yet in modern computer games, modern graphics cards denoise a scene in real-time at 60 frames per second using machine learning models [1][2] while doing all the other rendering at the same time. Granted, that's ray tracing, and the resolution is lower, and they technically cheat by using additional information, but it might be that DXO is not optimized very well. 1: https://blogs.nvidia.com/blog/ai-decoded-ray-reconstruction/ 2: https://gpuopen.com/amd-fsr-rayregeneration/

GDB supports it https://sourceware.org/gdb/current/onlinedocs/gdb.html/AMD-GPU.html You also get UMR from AMD https://gitlab.freedesktop.org/tomstdenis/umr There is also a bunch of other tools provided: https://gpuopen.com/radeon-gpu-detective/ https://gpuopen.com/news/introducing-radeon-developer-tool-suite/

It's essentially a replacement for vertex shaders which map more closely to how GPUs are processing big and complex triangle meshes as small packets of vertices in parallel by doing the job of splitting a complex triangle mesh into such small packets of vertices in an offline asset-pipeline job instead of relying too much on 'hardware magic' like vertex caches. AFAIK mesh shaders also get rid of (the ever troublesome) geometry shaders and hull shaders, but don't quote me on that :) By far most traditional triangle rendering use cases should only see minimal performance improvements though, it's very much the definition of 'diminishing returns'. It's definitely more straightforward and 'elegant' though. PS: this is a pretty good introduction I think https://gpuopen.com/learn/mesh_shaders/mesh_shaders-from_vertex_shader…

This used to be the case. However, they have published docs on their hardware. For instance, https://gpuopen.com/amd-gpu-architecture-programming-documentation/ Their driver is also in-tree: https://wiki.archlinux.org/title/AMDGPU It may or may not be as good as Intel. I don't have the depth of experience there right now. It looks like AMD and Intel GPUs should be roughly similar in terms of support / out of the box experience on Linux, but AMD will have better performance. Nvidia is still super proprietary, so it will generally have more headaches and hoops to deal with. So, generally, I rank them by default: 1. AMD 2. Intel 3. Nvidia Unless there is a reason (firmware or configuration differences for Intel/AMD needing to run/develop CUDA code for Nvidia) I try to stick with that. I've not yet had an AMD laptop though.

For AMD you can use the Radeon GPU Analyzer: https://gpuopen.com/rga/

It's all still available, arguably more accessible from a cost and documentation point of view - I'd argue that the bar to getting something is so much higher. You can today start banging registers on GPUs, or write your own shader programs, there's well documented modern hardware (e.g. [0]) and completely open-source drivers just sitting there for reference. The issue is that there's so moving parts now that to get anything visible to actually happen it's not just wiring up a couple of registers, but thousands just to get it into a state you can even see the registers you want, then a thousand more. [0] https://gpuopen.com/amd-gpu-architecture-programming-documentation/

Here is one article on this topic: https://gpuopen.com/learn/mesh_shaders/mesh_shaders-index/

GPUs have wavefronts so I assume it is similar? Here is a page that explains it: https://gpuopen.com/learn/occupancy-explained/

Ah, you watched my "good parallel computer" talk[﹡] :) If I were going to take it on, I'd start with BVH construction - the H-PLOC paper at the latest HPG [1] looks promising - then traverse down the hierarchy until you get very small number of path segments so you can pairwise compare them. Obviously any time there is an intersection you need at least the two segments. This seems hard to me, humbling even. I mean, overlap removal is hard enough on the CPU, especially because it's so sensitive to numerical robustness, and doubly so for curves. But I think you'll learn something for trying! [﹡] https://news.ycombinator.com/item?id=41105102 but it didn't make the front page; I'm holding back promoting it further pending writing a companion blog post. [1]: https://gpuopen.com/download/publications/HPLOC.pdf

Related: https://gpuopen.com/download/publications/Real-Time_Procedural_Generat… https://github.com/GPUOpen-LibrariesAndSDKs/WorkGraphComputeRasterizer https://gpuopen.com/fidelityfx-cauldron-framework/

Related: https://gpuopen.com/learn/work_graphs_mesh_nodes/work_graphs_mesh_node…

At the very least AMD made an extension pretty much immediately[1,2]. NVidia... well I don’t really expect graphics-programming-friendly initiatives from NVidia at this point. [1] https://gpuopen.com/gpu-work-graphs-in-vulkan/ [2] https://github.com/KhronosGroup/Vulkan-Docs/blob/main/proposals/VK_AMD…

This is quite crude on what it is. e.g. how does it differ from gpuopen[0] ISA documentation released by AMD themselves? 0. https://gpuopen.com/amd-isa-documentation/

preshaping before the 3d lut, such as with a 1d lut, is a common way to handle hdr tone mapping as well. See for example https://gpuopen.com/wp-content/uploads/2016/03/GdcVdrLottes.pdf

On the claim that instruction sets are undocumented: https://gpuopen.com/amd-isa-documentation/ Of course this doesn't apply to Nvidia, but MI-300 seems to be pretty viable for machine learning, so if you want things to change across the industry, more people need to put their money where their mouths are.

Not sure if I completely understand what "Ryzen AI" does, but tinygrad for example has some limited support for RDNA3[0]. It isn't quite there yet in matters of performance though, as you can read in the comments of that file. There's also a small tutorial by AMD on how to use the WMMA intrinsics[1] using their hipcc[2] compiler. Documentation is sparse kinda sparse, but the instruction set is not huge. The RDNA3 ISA guide[3] might also be helpful (and only a fraction of the pages are relevant.) So it doesn't seem like RE is necessary? AMD dumped most code on GitHub. 0. https://github.com/tinygrad/tinygrad/blob/master/extra/gemm/hip_matmul… 1. https://gpuopen.com/learn/wmma_on_rdna3/ 2. https://github.com/ROCm/HIPCC 3. https://www.amd.com/content/dam/amd/en/documents/radeon-tech-docs/inst…

I would guess the documentation is something of a hindrance. I'm in compilers rather than drivers, where the public docs are essentially the pdfs at https://gpuopen.com/amd-isa-documentation/. It's a lot better than the nvidia equivalent but it would also be rough going implementing a compiler based on those alone; too many instructions with names but unspecified semantics. The driver / firmware story is probably similar.

AMD isn't directly involved with the alliance, but they have embraced USD: https://www.amd.com/en/technologies/radeon-prorender ("A Pixar® USD™ Hydra™ render delegate...") https://gpuopen.com/learn/amd-usd-hydra-blender