3 ms·

> but its still new/delete based manual memory allocation I guess we aren’t solving the same problems: memory allocation is trivial is my domain, mapping compl

> but its still new/delete based manual memory allocation I guess we aren’t solving the same problems: memory allocation is trivial is my domain, mapping complex nested control flow in SIMD is the hard part. > Julia can now compile down into PTX or AMDGPU assembly. Sure but Julia-as-syntax is nothing special now, Numba does this for a Python as well.

dnautics 6y ago
you can and you can't. Julia is composable -- Suppose I want to write a library to find compression polynomials for reed-solomon encoding system (see mary wootters' talks on youtube) for my storage product. I need a LU decomposition algorithm that operates on GF256 (which is just an 8-bit int). But the +/-/x/divide operations are all messed up. I'd have to rewrite LU decomposition. How confident are you that you can get even that right? I'm pretty good at implementing algorithms, but I've messed up LU decomposition before. Then suppose I rewrite the LU decomposition algorithm, using python. Now I want to accelerate the search by running the search on GPUs. I have to re-rewrite the code from scratch. Each GF256 encoding has to have rejiggered operators, and so I need to rewrite custom GPU kernels, then figure out how to resequence the operations (* looks different for each GF256 encoding), etc. This is all SUPER easy in julia.
- marmaduke 6y ago
  Then use Julia, that's great. My comment was that the SIMD support in Julia isn't sufficient for my problem, not that you can't do your GF256 linear alg on a GPU..?
- fnord123 6y ago
  > (see mary wootters' talks on youtube) https://www.youtube.com/watch?v=Gh578e98qAk https://www.youtube.com/watch?v=Gh578e98qAk > Suppose I want to write a library to find compression polynomials for reed-solomon encoding system (see mary wootters' talks on youtube) for my storage product. I need a LU decomposition algorithm that operates on GF256 Surely you use isa-l[1]. [1] https://github.com/intel/isa-l https://github.com/intel/isa-l > Now I want to accelerate the search by running the search on GPUs. GPUs are float oriented so I don't think you'll get the performance you hope for out of 8 bit integer operations. If you have interesting results to share I'd like to read them.
  - dragontamer 6y ago
    > GPUs are float oriented so I don't think you'll get the performance you hope for out of 8 bit integer operations. If you have interesting results to share I'd like to read them. You've never seen GPU Hashcat, or GPU Bitcoin / Ethereum mining? Vega now has a huge focus on INT8 operations. NVidia can perform int operations in parallel with float operations (superscalar GPU cores)
    - fnord123 6y ago
      > You've never seen GPU Hashcat, or GPU Bitcoin / Ethereum mining? I've heard of it, but afaik it hasn't been profitable to mine using a GPU for a long time due to the competition for hashes, power consumption, and rate they mine at. This is in contrast to ASICs which can mine even faster for less capex and opex.
dragontamer 6y ago
> I guess we aren’t solving the same problems: memory allocation is trivial is my domain, mapping complex nested control flow in SIMD is the hard part. Those statements are just confusing to me. In most GPU-code, you have crazy amounts of parallelism and therefore don't really care what order those statements execute in. As such, if you can allocate memory, and map those if/else statements into a collection of queues and/or stacks (depending on if you want a breadth-first search, or depth-first search pattern). In effect: you use memory allocation to solve the complex control flow issue. Maybe its more obvious with code: Instead of doing: if(baz()){ foo(); } else { bar(); // Thread divergence!! } Do: if(baz()){ pushIntoFooQueue(); } else { pushIntoBarQueue(); // Thread divergence, but not much of a penalty } while(fooIsNotEmpty()){ // No thread divergence at all task = SIMDPopFoo(); task.execute(); } while(barIsNotEmpty()){ task = SIMDPopBar(); task.execute(); } This is heavier on the memory units, because you now have to manage the data-structures. But this style practically negates the thread-divergence penalty completely. If you're lucky, your fooQueue and barQueue fit in __shared__ memory. Bonus points: Not only is thread-divergence negated, but you also achieve effective load-balancing across your workgroup. If Thread#0 spawns 20 items for FooQueue, after pushing/popping from the queue, those 20 Foo-tasks will be assigned to 20 different threads. ----- From there on out, you're just pushing / popping different parts of your code to various queues and/or stacks. But this is only really a valid solution if you have a decent memory allocator that knows where and how to clean up these queues / stacks (especially if you have nodes starting to point to each other for dependency management) I haven't really solved this problem "in general", but is clear that the queues should be sorted into topological order, and that any tasks that depend onto each other need to be run in different iterations. It really depends on how much you're willing to spend on organizing this execution information. In any case: the memory allocation issue is one-and-the-same with the thread-divergence / complex instruction flow issue to me. You need to create a data-structure that organizes the instruction flow, and any complex data-structure will need memory management as soon as you start linking things together. The above uses a queue or stack, but things can get more complicated. -------- Not that Julia, Python, ISPC or anything really... solves this problem. But memory management is very useful for this "style" (I'm mostly doing ref-counted C++ with a custom allocator myself. But such trees or graphs of links can grow into the CPU-side and end up using the default memory allocator)

dnautics 6y ago

you can and you can't. Julia is composable -- Suppose I want to write a library to find compression polynomials for reed-solomon encoding system (see mary wootters' talks on youtube) for my storage product. I need a LU decomposition algorithm that operates on GF256 (which is just an 8-bit int). But the +/-/x/divide operations are all messed up. I'd have to rewrite LU decomposition. How confident are you that you can get even that right? I'm pretty good at implementing algorithms, but I've messed up LU decomposition before. Then suppose I rewrite the LU decomposition algorithm, using python. Now I want to accelerate the search by running the search on GPUs. I have to re-rewrite the code from scratch. Each GF256 encoding has to have rejiggered operators, and so I need to rewrite custom GPU kernels, then figure out how to resequence the operations (* looks different for each GF256 encoding), etc. This is all SUPER easy in julia.

marmaduke 6y ago
Then use Julia, that's great. My comment was that the SIMD support in Julia isn't sufficient for my problem, not that you can't do your GF256 linear alg on a GPU..?
fnord123 6y ago
> (see mary wootters' talks on youtube) https://www.youtube.com/watch?v=Gh578e98qAk https://www.youtube.com/watch?v=Gh578e98qAk > Suppose I want to write a library to find compression polynomials for reed-solomon encoding system (see mary wootters' talks on youtube) for my storage product. I need a LU decomposition algorithm that operates on GF256 Surely you use isa-l[1]. [1] https://github.com/intel/isa-l https://github.com/intel/isa-l > Now I want to accelerate the search by running the search on GPUs. GPUs are float oriented so I don't think you'll get the performance you hope for out of 8 bit integer operations. If you have interesting results to share I'd like to read them.
- dragontamer 6y ago
  > GPUs are float oriented so I don't think you'll get the performance you hope for out of 8 bit integer operations. If you have interesting results to share I'd like to read them. You've never seen GPU Hashcat, or GPU Bitcoin / Ethereum mining? Vega now has a huge focus on INT8 operations. NVidia can perform int operations in parallel with float operations (superscalar GPU cores)
  - fnord123 6y ago
    > You've never seen GPU Hashcat, or GPU Bitcoin / Ethereum mining? I've heard of it, but afaik it hasn't been profitable to mine using a GPU for a long time due to the competition for hashes, power consumption, and rate they mine at. This is in contrast to ASICs which can mine even faster for less capex and opex.

dragontamer 6y ago

> I guess we aren’t solving the same problems: memory allocation is trivial is my domain, mapping complex nested control flow in SIMD is the hard part. Those statements are just confusing to me. In most GPU-code, you have crazy amounts of parallelism and therefore don't really care what order those statements execute in. As such, if you can allocate memory, and map those if/else statements into a collection of queues and/or stacks (depending on if you want a breadth-first search, or depth-first search pattern). In effect: you use memory allocation to solve the complex control flow issue. Maybe its more obvious with code: Instead of doing: if(baz()){ foo(); } else { bar(); // Thread divergence!! } Do: if(baz()){ pushIntoFooQueue(); } else { pushIntoBarQueue(); // Thread divergence, but not much of a penalty } while(fooIsNotEmpty()){ // No thread divergence at all task = SIMDPopFoo(); task.execute(); } while(barIsNotEmpty()){ task = SIMDPopBar(); task.execute(); } This is heavier on the memory units, because you now have to manage the data-structures. But this style practically negates the thread-divergence penalty completely. If you're lucky, your fooQueue and barQueue fit in __shared__ memory. Bonus points: Not only is thread-divergence negated, but you also achieve effective load-balancing across your workgroup. If Thread#0 spawns 20 items for FooQueue, after pushing/popping from the queue, those 20 Foo-tasks will be assigned to 20 different threads. ----- From there on out, you're just pushing / popping different parts of your code to various queues and/or stacks. But this is only really a valid solution if you have a decent memory allocator that knows where and how to clean up these queues / stacks (especially if you have nodes starting to point to each other for dependency management) I haven't really solved this problem "in general", but is clear that the queues should be sorted into topological order, and that any tasks that depend onto each other need to be run in different iterations. It really depends on how much you're willing to spend on organizing this execution information. In any case: the memory allocation issue is one-and-the-same with the thread-divergence / complex instruction flow issue to me. You need to create a data-structure that organizes the instruction flow, and any complex data-structure will need memory management as soon as you start linking things together. The above uses a queue or stack, but things can get more complicated. -------- Not that Julia, Python, ISPC or anything really... solves this problem. But memory management is very useful for this "style" (I'm mostly doing ref-counted C++ with a custom allocator myself. But such trees or graphs of links can grow into the CPU-side and end up using the default memory allocator)