3 ms·

Unfortunately, I don't see a "just a bit of magic here without learning much of anything new" interface coming because it's all about strategizing the movement

Unfortunately, I don't see a "just a bit of magic here without learning much of anything new" interface coming because it's all about strategizing the movement of data. This is not unique to GPUs. It's a universal problem across computing hardware. It's just enabled to be explicit in OpenCL/CUDA. As compared to most languages where you indirectly try to steer things the right way and the CPU does it's best with whatever mess it gets. Closest I know of is https://halide-lang.org/ https://halide-lang.org/ And, that is specialized around images.

dragontamer 5y ago
Both C++AMP and Tensorflow abstracted away the movement of data behind relatively easy to use interfaces. C++ AMP in particular was a generic interface behind a template<> class: Array<> and ArrayView<>. Array<> always existed on GPU-side (be it GPU#0 or GPU#1), while ArrayView<> was your abstraction to read/write into that Array<>. Array<SomeClass> foo = createSomeArray(); ArrayView<SomeClass> bar = foo; This could be CPU-side or GPU side code (foo could be on GPU#0, while bar could be GPU#1). While you can use C++ AMP code to manually manage when the copies happen, the idea is to simplify the logic down so that CPU-programmers didn't have to think as hard. In this case, "bar" is bound to foo, and bar will remain in sync with foo through async programming behind the scenes. bar[25] = SomeClass(baz); will change bar in this context, but the Array<> foo will only update at the designated synchronization points (similar to CUDA Streams). ------ So its a lot easier than reading/writing cudaMemcpy() in the right spots, in the right order, in the right CUDA Streams. But probably less efficient, and therefore going to be less popular than using CUDA directly.
- my123 5y ago
  > and therefore going to be less popular than using CUDA directly. The current ambition of NV is trying to make it as painless as possible, by relying on shared virtual memory much more than before. https://www.olcf.ornl.gov/wp-content/uploads/2021/06/OLCF_User_Call_July_2021_HPC-SDK.pdf https://www.olcf.ornl.gov/wp-content/uploads/2021/06/OLCF_Us... With OpenMP (also supported in AMD’s ROCm AOMP) and C++17 standard parallelism support. Legate w/ cuNumeric goes to another direction with implementing a numpy-compatible interface, but GPU accelerated. Also deriving from the limits of a SIMD abstraction with 1 IP per thread for the sake of guaranteed forward progress (as such changing the programming model needed for the code to be functional). The underlying machine is implemented as SIMT, so people should still be aware of the associated performance pitfalls. However this allows for better compatibility and more code sharing possible than before. Microsoft’s problem with C++ AMP is that FXIL as present in D3D11 was just too limited and had too many performance pitfalls I’m afraid, if only it got to see the light of day with D3D12 + DXIL later…
  - dragontamer 5y ago
    > The current ambition of NV is trying to make it as painless as possible, by relying on shared virtual memory much more than before. Shared-virtual-memory is an OpenCL term. I think the NVidia term is "Unified Memory", but yes, I get what you're trying to say. I think that makes the mistake too far in the simplicity direction. Its over-simplified. Like, think about the following code: DataStructure data = foo(); // cudaMallocManaged() / Shared-virtual Memory bar(data); baz(data); In a shared-virtual memory situation, you'd probably have data sitting on CPU-side, and then SVM over the data around. With a little bit more code, you can have... Array<DataStructure> data = foo(); ArrayView<DataStructure> dataview = data; bar(data); baz(data); In this case, the data exists on GPU-side (explicitly: all Array<> in C++Amp are GPU-allocated). If foo, bar and baz are all GPU-functions, then the data never moves (no cudaMemcpy involved). If bar is a CPU-function, you get the copy back-and-forth for correctness. But that still leaves room for you to convert bar from CPU-code into GPU-code without affecting this section of code... giving you the ability to smoothly improve performance as your project matures. ------- CUDA / OpenCL default memory is too granular. Datastructure* data = cudaMalloc(...); // GPU allocation Datastructure* localData = malloc(...); cudaMemcpy(localData, data, stream-stuff, etc. etc.) // ugggh... too much. Lots of steps to do anything, thinking about streams and async and signals. ------- Finding the right compromise between performance and simplicity is key. CUDA/OpenCL have both something that's too simple(SVM / Unified Memory), and too complicated (cudaMalloc + cudaMemcpy). The answer IMO lies in-between somehow.
    - my123 5y ago
      > CUDA/OpenCL have both something that's too simple(SVM / Unified Memory) It’s sadly the only choice around for maximum code reuse… with then manually optimising the parts that need to be later through profiling, as developer time is quite limited. On systems where the CPU and GPU are coherent such SoCs w/ CPU and GPU on the same die, Infinity Fabric-linked GPUs or NVlink linked ones (Power9/Grace), that option is also quite attractive. The memory manager measures the frequency of access from the CPU and GPU to (try to) locate where a buffer should be stored or moved dynamically. However, on PCIe linked GPUs, reliance on heuristics is more necessary to reach good perf levels. You can also ask to have your memory allocation migrated to the GPU on-demand with cudaMemPrefetchAsync (finally implemented in ROCm 4.5 too, which has unified memory on AMD cards) before doing a kernel launch for the data that you’ll use. (See https://on-demand.gputechconf.com/gtc/2017/presentation/s7285-nikolay-sakharnykh-unified-memory-on-pascal-and-volta.pdf https://on-demand.gputechconf.com/gtc/2017/presentation/s728... for the changes that happened in Volta on that front) The goal is to make the use of regular cudaMalloc/hipMalloc not necessary for good performance in a wide variety of scenarios. The explicit storage location APIs will stay available for when more optimisation is necessary.
convolvatron 5y ago
there were the '' languages on the connection machine. in C for example, 'shapes', which were cartesian arrays of values were part of the type system. communications operations were explicit through the use of indirection on these indices. I doubt its the case on a modern gpu, but aside from the router, performance on the CM was deterministic so you really could understand what you were going to get from the source. completely agree with you, general parallel programming languages are a thing and really pretty valuable to try to exploit all that machinery. Chapel is the only recent work I'm familiar with.
- dragontamer 5y ago
  You got "italics" because * turns into italics in this website. I'm pretty sure you are talking about Lisp* and C*, which were the "star" languages of the connection machine.
joe_the_user 5y ago
I think the thing is that the only reason to do general purpose computing on a GPU is to get extra performance. So everyone who's using the device is going to want to know how to optimize where a program where performance doesn't matter will default run on a CPU. And GPUs are structured to send a whole lot of memory to many processors and I think that doing this efficiently is an inherently harder problem than sending memory to a few complex processors through a pipeline that's complex, yes, but designed to make things feel easy.

dragontamer 5y ago

Both C++AMP and Tensorflow abstracted away the movement of data behind relatively easy to use interfaces. C++ AMP in particular was a generic interface behind a template<> class: Array<> and ArrayView<>. Array<> always existed on GPU-side (be it GPU#0 or GPU#1), while ArrayView<> was your abstraction to read/write into that Array<>. Array<SomeClass> foo = createSomeArray(); ArrayView<SomeClass> bar = foo; This could be CPU-side or GPU side code (foo could be on GPU#0, while bar could be GPU#1). While you can use C++ AMP code to manually manage when the copies happen, the idea is to simplify the logic down so that CPU-programmers didn't have to think as hard. In this case, "bar" is bound to foo, and bar will remain in sync with foo through async programming behind the scenes. bar[25] = SomeClass(baz); will change bar in this context, but the Array<> foo will only update at the designated synchronization points (similar to CUDA Streams). ------ So its a lot easier than reading/writing cudaMemcpy() in the right spots, in the right order, in the right CUDA Streams. But probably less efficient, and therefore going to be less popular than using CUDA directly.

my123 5y ago
> and therefore going to be less popular than using CUDA directly. The current ambition of NV is trying to make it as painless as possible, by relying on shared virtual memory much more than before. https://www.olcf.ornl.gov/wp-content/uploads/2021/06/OLCF_User_Call_July_2021_HPC-SDK.pdf https://www.olcf.ornl.gov/wp-content/uploads/2021/06/OLCF_Us... With OpenMP (also supported in AMD’s ROCm AOMP) and C++17 standard parallelism support. Legate w/ cuNumeric goes to another direction with implementing a numpy-compatible interface, but GPU accelerated. Also deriving from the limits of a SIMD abstraction with 1 IP per thread for the sake of guaranteed forward progress (as such changing the programming model needed for the code to be functional). The underlying machine is implemented as SIMT, so people should still be aware of the associated performance pitfalls. However this allows for better compatibility and more code sharing possible than before. Microsoft’s problem with C++ AMP is that FXIL as present in D3D11 was just too limited and had too many performance pitfalls I’m afraid, if only it got to see the light of day with D3D12 + DXIL later…
- dragontamer 5y ago
  > The current ambition of NV is trying to make it as painless as possible, by relying on shared virtual memory much more than before. Shared-virtual-memory is an OpenCL term. I think the NVidia term is "Unified Memory", but yes, I get what you're trying to say. I think that makes the mistake too far in the simplicity direction. Its over-simplified. Like, think about the following code: DataStructure data = foo(); // cudaMallocManaged() / Shared-virtual Memory bar(data); baz(data); In a shared-virtual memory situation, you'd probably have data sitting on CPU-side, and then SVM over the data around. With a little bit more code, you can have... Array<DataStructure> data = foo(); ArrayView<DataStructure> dataview = data; bar(data); baz(data); In this case, the data exists on GPU-side (explicitly: all Array<> in C++Amp are GPU-allocated). If foo, bar and baz are all GPU-functions, then the data never moves (no cudaMemcpy involved). If bar is a CPU-function, you get the copy back-and-forth for correctness. But that still leaves room for you to convert bar from CPU-code into GPU-code without affecting this section of code... giving you the ability to smoothly improve performance as your project matures. ------- CUDA / OpenCL default memory is too granular. Datastructure* data = cudaMalloc(...); // GPU allocation Datastructure* localData = malloc(...); cudaMemcpy(localData, data, stream-stuff, etc. etc.) // ugggh... too much. Lots of steps to do anything, thinking about streams and async and signals. ------- Finding the right compromise between performance and simplicity is key. CUDA/OpenCL have both something that's too simple(SVM / Unified Memory), and too complicated (cudaMalloc + cudaMemcpy). The answer IMO lies in-between somehow.
  - my123 5y ago
    > CUDA/OpenCL have both something that's too simple(SVM / Unified Memory) It’s sadly the only choice around for maximum code reuse… with then manually optimising the parts that need to be later through profiling, as developer time is quite limited. On systems where the CPU and GPU are coherent such SoCs w/ CPU and GPU on the same die, Infinity Fabric-linked GPUs or NVlink linked ones (Power9/Grace), that option is also quite attractive. The memory manager measures the frequency of access from the CPU and GPU to (try to) locate where a buffer should be stored or moved dynamically. However, on PCIe linked GPUs, reliance on heuristics is more necessary to reach good perf levels. You can also ask to have your memory allocation migrated to the GPU on-demand with cudaMemPrefetchAsync (finally implemented in ROCm 4.5 too, which has unified memory on AMD cards) before doing a kernel launch for the data that you’ll use. (See https://on-demand.gputechconf.com/gtc/2017/presentation/s7285-nikolay-sakharnykh-unified-memory-on-pascal-and-volta.pdf https://on-demand.gputechconf.com/gtc/2017/presentation/s728... for the changes that happened in Volta on that front) The goal is to make the use of regular cudaMalloc/hipMalloc not necessary for good performance in a wide variety of scenarios. The explicit storage location APIs will stay available for when more optimisation is necessary.

convolvatron 5y ago

there were the '' languages on the connection machine. in C for example, 'shapes', which were cartesian arrays of values were part of the type system. communications operations were explicit through the use of indirection on these indices. I doubt its the case on a modern gpu, but aside from the router, performance on the CM was deterministic so you really could understand what you were going to get from the source. completely agree with you, general parallel programming languages are a thing and really pretty valuable to try to exploit all that machinery. Chapel is the only recent work I'm familiar with.

dragontamer 5y ago
You got "italics" because * turns into italics in this website. I'm pretty sure you are talking about Lisp* and C*, which were the "star" languages of the connection machine.

joe_the_user 5y ago

I think the thing is that the only reason to do general purpose computing on a GPU is to get extra performance. So everyone who's using the device is going to want to know how to optimize where a program where performance doesn't matter will default run on a CPU. And GPUs are structured to send a whole lot of memory to many processors and I think that doing this efficiently is an inherently harder problem than sending memory to a few complex processors through a pipeline that's complex, yes, but designed to make things feel easy.