GPUCC – An Open-Source GPGPU Compiler
research.google.com
research.google.com
It says that a lot of the reason TensorFlow initially lagged in performance is because a lot of those performance issues only manifested under NVCC, whereas they had been using GPUCC internally.
Can anyone comment on the following quote:
The list below shows some of the more important optimizations for GPUs... A few of them have not been upstreamed due to lack of a customizable target-independent optimization pipeline.
So the LLVM version of gpucc will be incomplete? Will there be a release of the original stand-alone gpucc?
Yes, it is currently incomplete, but I'd say at least 80% of the optimizations are upstreamed already. Also, folks in the LLVM community are actively working on that. For example, Justin Lebar recently pushed http://reviews.llvm.org/D18626 that added the speculative execution pass to -O3.
Regarding performance, one thing worth noting is that missing one optimization does not necessarily cause significant slowdown on the benchmarks you care about. For example, the memory-space alias analysis only noticeably affects one benchmark in the Rodinia benchmark suite.
Regarding your second question, the short answer is no. The Clang/LLVM version uses a different architecture (as mentioned in http://wujingyue.com/docs/gpucc-talk.pdf) from the internal version. The LLVM version offers better functionality and compilation time, and is much easier to maintain and improve in the future. It would cost even more effort to upstream the internal version than to make all optimizations work with the new architecture.
I'm definitely subscribing to the llvm-dev list[1] in case any discussion on this continues there. There's also the llvm-commits, clang-dev, and clang-commits lists as well, but llvm-dev kinda seems like the right place for this.
Gpucc in LLVM is definitely a breath of fresh air for all of us nvcc users. To get to see some compiler internals for cuda, it feels like Christmas. A big thanks from me for all the upstreaming effort!
I don't have a lot of benchmarks at the moment, so I can't say how important they are. And it of course depends on what you're doing.
clang/llvm's CUDA implementation shares most of the backend with gpucc, but it's an entirely new front-end. The front-end works for tensorflow, eigen, and thrust, but I suspect if you try hard enough you'll be able to find something nvcc accepts that we can't compile. At the moment we're pretty focused on making it work well for Tensorflow.
http://www.pgroup.com/doc/pgicudaforug.pdf
One could take one's Fortran code and simply recompile it with their compiler to run on the Nvidia GPU's. The compiler would perform automatic parallelization. Wild stuff.
Btw. I'm working on something that's geared towards pretty much exactly what you're talking about. My stretch goal is fully automatic GPU parallelization for data parallel Fortran code [1].
I don't see anything specifically about this in the archives, but maybe that's something to look forwards to.
When I look at CUDA code, it seems to be a big loop targeting the GPU memory with standard c code, allocating memory with standard functions and specifying where code lives with simple defines.
When I look at OpenCL, it is... I don't know what it is. I haven't figure it out after considerable scanning. And that has cemented my decision to avoid it because I don't have infinite time to scan obscurity.
For example, here is a standard "first OpenCL program" - ~200 lines of boiler plate and no simple example of our many cores working together to do something brutally simple and useful like add two vectors. Just "hello world" from GPU.
As far as I can tell, as a production of a multitude of vendors all of which have different stuff, OpenCl is a monstrosity where you have a wide of variety of functionalities supported but none of those functionalities is guaranteed to be present - hence 200 lines of boiler plate. Kind of like the umpteen Unix flavors and such back in the day, "Open standards" that are bridges between only semi-compatible hardware have generally been doomed abortions discarded in favor of a single best approach that all vendors are forced to adopt.
So it seems like the best thing is jettisoning the monstrosity and cloning CUDA for other hardware.
https://www.fixstars.com/en/opencl/book/OpenCLProgrammingBoo...
Use the cl C++ wrapper if brevity is important to you and using C++ is a choice. The hello world example here is noticeably shorter
http://simpleopencl.blogspot.com/2013/06/tutorial-simple-sta...
The problem is that the "canonical example" pretty much remains what I showed. And what's bad about that example isn't simple length but the way that the creation and manipulation of kernels and threads remains entirely opaque (in contrast to your example I think).
What you are looking at is largely just code to compile a kernel and launch it on the device, basically a makefile in C. You can really just ignore most of the boilerplate until you need to modify it.
"Hello World" is not a particularly meaningful program on a device with no console or LEDs to blink - why don't you try something simple like summing every number from 1 to N? Or, a Monte-Carlo simulator to calculate Pi? Easy to code, easy to verify...
CUDA is also quite verbose if you are playing by all the rules. You are still supposed to specify your target device, verify your capabilities, free your buffers, check your async return values, etc - and if you don't then you will get weird behavior down the road until you figure out what's going on. OpenCL might try to helpfully ignore misbehavior too, I don't know for sure.
Personally I find both CUDA and OpenCL to be quite verbose and when I write for them I use a library like Thrust or Bolt, which takes most of the boilerplate out of it and feel much like writing C++ STL code. It also allows you to work efficiently at a much higher level, and then come through after-the-fact and optimize the parts that are actually slowing you down. It also provides stuff like automatic occupancy tuning and so on. You can trivially switch between Thrust and native kernels by using raw_pointer, so it's great for sorts, scans, etc that "glue" things together. Use it like you would Python, bearing in mind that round trips to host memory are slow (but perhaps not fatally so!)
http://docs.nvidia.com/cuda/thrust/
http://developer.amd.com/tools-and-sdks/opencl-zone/bolt-c-t...
One useful trick is that you can write your CUDA functions inside a Thrust functor which iterates your data elements. So you can use a count_iterator which represents the index of the data element(s) or grid processing element in question, and then you write a functor which uses the counter value to load a data element from a pointer stored inside the functor, and does some work on it. This gives you a "kernel simulator" as an intermediate step between array-wide functional programming and native __device__ functions.
This also provides a great place for indirection so you can easily swap between a GPU backend and Thrust's OpenMP CPU backend.
I'm not trying to run "hello world", I'm looking at the easily accessible OpenCl code and seeing "hello world" as the most common example program.
Any "easy wrappers" for c?
I think Khronos wrapper is just a translation of the original API into c++, which partakes of the basic problem - a zillion options for a zillion distinct sorts of hardware, in contrast to Cuda, with a main functional and supported approach.
See, for, example, the streamexecutor thread, etc.
Also, outside of that, they've been upstreaming the non-controversial smaller stuff that is part of this for many months now.
(Seriously, i think of all the companies you are going to complain about, you may want to look at google's interactions with clang and llvm, where we are actually one of the only folks who work completely upstream at all times, before throwing stones)
But I know some things about what words mean, and publishing a paper describing an open source project in March and not having any code available for download in April is just kind of weird, no?
It's good to talk about working out the precise mechanics of upstreaming code. But in an open source project, you'd expect to publish your fork so that other people can play an informed part in that conversation.
In an open source project, the best expectation is to play by the rules of the existing project and try to integrate your changes in the way that works well with it. There's no One True Open Source way - there are a lot of projects, each with their own cultures.
What Google's done is take an internally developed thingy and transition it to LLVM. That's a pretty non-trivial effort for any company. I don't see why having it be open source has any requirement for a dump of the internal version. Earlier commenters noted already that a large fraction of the code is already present in LLVM and has been streaming in for some time now, so -- why the hostility?
Collectively, I don't think "our" (the wider community) goal is necessarily to have a bunch of junk forks out there that can't be compiled or used. Working, thoughtfully contributed code is much more likely to be widely used and have a big impact, and that's a standard we should be happy if companies meet. LLVM isn't a quick hack project - it's a foundational bit of tech that millions of people depend on directly or indirectly, and that millions of people benefit from improvements to.
This is a lot of assumptions. For starters, it's all pretty much part of llvm 3.9 already.
"It's good to talk about working out the precise mechanics of upstreaming code. But in an open source project, you'd expect to publish your fork so that other people can play an informed part in that conversation. "
Actually, no, you wouldn't, at least not in these communities. In fact, in most of the projects i've been in, that's the last thing i'd want or the community would want, because it encourages people to play with these weird hybrid versions, when that's not what anyone really wants. Instead, we'd want people to come to us with design ideas, and use cases, not a fait accompli already written that they expect us to do something with.
I'm not sure what projects you are part of that work the opposite way, where it's "fork everything and whatever happens happens", but it's not the majority of projects i've belonged to nor have i personally found it to be a very valuable mechanism. It makes people attached to their current implementations. When it comes to things like this, a lot of them are written the way they are not because it's a good design, but because they needed to get stuff done. Thus, discussing the use case, and design, and providing example code if people want it to play with, great. Dumping forks on the world, not a particularly useful thing in most cases except for boosting egos, IMHO (there are times it is useful,for sure, but those are mostly rare cases IMHO).
Look at the history of GCC and LLVM forks. Of those who published forks and started discussion about merging those forks as-is, the number of successful merges is near zero.
Of those who said 'hey, we did some cool stuff, here's our use cases and our initial thoughts on design. We have stuff that follows this design we can show but are otherwise happy to figure out the right design and build that", the number of successful merges is near 100%.
So that's what was done.
http://gpuopen.com/compute-product/hcc-heterogeneous-compute... https://github.com/RadeonOpenCompute/hcc
It would take all of getting to page 2 of the article to confirm this instead of speculating...
OTOH, there is an intriguing footnote that
> We are also experimenting compiling [virtual ISA] PTX to [Nvidia's proprietary Shader ASSembler] SASS before program execution and embedding the SASS directly into the resultant binary
but the paper mentions in the conclusion that a SASS spec is not publicly available. It would be interesting for someone involved to comment more on that. Experiments on reverse engineering the compiled PTX results?
If implementing a replacement for nvcc gave these gains, I would imagine being able to control an offline version of the (normally JIT) compilation to SASS would also yield large benefits. It would likely be incredibly architecture dependent, but for the big machine learning projects that still might be worth the expense.
I'm one of the people implementing this stuff in upstream clang/llvm.
I'm not aware of any experiments internal or otherwise where we've tried to compile to SASS. We believe that there would be some performance gains to be had, but it's a really big project, even if we had an ISA, which at the moment we do not.