Adding advanced FAQ #47
Conversation
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
| ``` | ||
| instead of pre-declaring and using `var[...]`. | ||
|
|
||
| ### When should I use `+=` operator? |
There was a problem hiding this comment.
This can be removed. +=, /=, -=, and *=n, works across sbuf and psum now. It has not been released, but the issue is fixed and folks can go back to using this syntax.
| Any other use of `+=` may trigger compiler bugs. Please do not use `+=` outside of matrix multiplication. | ||
|
|
||
| ### What's the difference between `affine_range` and `sequential_range`? | ||
| - [`nl.affine_range`](https://awsdocs-neuron.readthedocs-hosted.com/en/latest/general/nki/api/generated/nki.language.affine_range.html#nki.language.affine_range) creates a sequence of numbers for use as parallel loop iterators in NKI. What this means is that the compiler will unroll your loop, identify the parts that can be parallelized, and then run them in parallel. For this reason, `affine_range` should be the default loop iterator choice when there is no loop-carried dependency. What that means is that you can should only use `nl.affine_range` when there is no depdency between steps in your loops, ie when the steps can be executed at any point in time. Please note that associative reductions are not considered loop carried dependencies in this context. A concrete example of associative reduction is multiple `nl.matmul` or `nisa.nc_matmul` calls accumulating into the same output buffer defined outside of this loop level. `nl.affine_range` allows parallel execution of loop iterations when there are no dependencies |
There was a problem hiding this comment.
At the moment, we should encourage everyone to use built-in range.
Nit on the affine_range definition.
Maybe integrate something like:
"When the compiler unrolls a loop specified by nl.affine_range, each iteration is treated as an independent slice that can be run concurrently without interacting with any other iterations. This type of loop works well for automatically writing simple vectorized loops when it does not contain loop-carried dependencies."
"When the compiler unrolls a loop specifed by nl.sequential_range, each iteration is unrolled in order. In addition to supporting loop-carried dependencies, this loop type works well for manually writing vectorized loops."
There was a problem hiding this comment.
or suggest affine_range is like open mp parallel_for
There was a problem hiding this comment.
Hi team, I would like to ask whether the internal code of affine_range runs in parallel on the two cores of trn1. Compared to nl.sequential_range, does it introduce additional communication overhead?
There was a problem hiding this comment.
Hey Dinghong! I think it's more parallelizing execution across available engines, not necessary across cores (but am curious to know if others disagree).
Generally affine range should introduce much better perf due to parallelizing.
The static range is just for getting started and debugging.
|
|
||
| ``` | ||
|
|
||
| Both options can work. But option 2 will be more efficient, as it will turn the `nl.sequential_range()` back to `nl.affine_range()` and get better throughput since the loop iterations don't need to wait for a shared chunk of memory to be updated. Tensor `w_temp` is not used outside of the loop so allocating it outside of the loop just adds an unnecessary loop-carried dependency to figure out. |
| **Solution:** | ||
| - Check tensor dimensions against hardware limits | ||
| - Verify tile sizes are within SBUF partition size | ||
| - Consider breaking large tensors into smaller tiles |
There was a problem hiding this comment.
- If you are using masks, be sure you're also propagating them across all relevant compute operations.
| **Common causes:** | ||
| - Using very small tiles for DMA operations | ||
| - Complex nested loops with small tile sizes | ||
| - Excessive unrolling |
There was a problem hiding this comment.
static_range will also cause this.
There was a problem hiding this comment.
instead of saying "Excessive unrolling", we should say huge tripcount in nki loops?
Issue #, if available:
Adding an advanced FAQ quide
Description of changes:
Description of common issues and how to resolve them
Testing:
Please see detailed unit test requirements in the CONTRIBUTING.md
nki.baremetalnki.benchmarkPull Request Checklist