Well, IMO it's much harder to take legacy code and port it to GPU. I see lots of tasks where people take like one or two parts of the problem, push it to GPU, do a few operations, and pull it back.

Frankly I think that's the wrong approach to begin with - you don't get good speedups that way. Pushing everything to-and-fro across a PCIe bus that is less than half the speed of DDR3 let alone DDR4 is not a recipe for success. Literally the only place where that's even successful is when you can do a sort-and-search or something similar that the GPU is really super good at.

You really need to be doing almost everything in VRAM as much as possible, and really carefully picking what goes across the bus, because that will bottleneck you, no question. And the problem is that a lot of legacy code is not written with any of these ideas in mind. They're not memory efficient.

I originally inherited legacy C-code that was at least third-hand (and the reason the prof wanted help was because nothing worked right), and took about a year of part time work to reverse-engineer it into a new C implementation that was actually workable, then thread it with OpenMP. The GPU conversion was year 2-3 of this project.

I'm certainly not going to say the reference/OpenMP implementation was a masterwork, and I didn't squeeze it for every drop of memory or performance. But I have zero question that the CUDA implementation was much better. From what I remember it consumed at most half the memory if not less, and was easier to scale up with more processor resources. The functional-esque style with structure-of-arrays worked really really well for that and I actually ended up backporting some features like the "sort-and-search" approach that helped speed the OpenMP implementation up somewhat too (wasn't huge but it was some).

Side note, the Thrust library can target OpenMP as a __device__ back end. So if you write using the Functor-style I outlined, you can write Thrust programs and run them on your CPU for debug/etc. That was another reason I went that route that I didn't really get a chance to explore.

Anyway, what I'm saying here is that from what I've seen, the approach of trying to plug GPUs into a key part of a complex legacy app is doomed to fail. You get like 1-3x speedup at most, often a slowdown. This is embedded programming, you need to boil your problem down to the absolute minimum possible problem, squeeze it as small as possible to maximize your VRAM (problem size), keep everything on the GPU and do as much processing as possible, and minimize your transfers over your bottlenecks. When you do your transfers - do them in bulk instead of one at a time.

It's a very different model from "strong" cores like a CPU, and you have to factor in that it's across a pretty slow bus (APUs with cache-coherent busses are a promising model, as is Knight's Landing). Offloading stuff to a co-processor isn't trivial to begin with, let alone when it has a weird programming model like a GPU that's very different from "strong" CPU cores.

Others have said that too, I really will try to clean this up and repost it. It'll probably end up being a series because a full explanation of each of those chunks will be a couple pages.