[NPU]: Add NPU support for the embedding #1028
Open
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Contributor
Author
|
test_embedding result： |
Contributor
Author
|
benchmark_embedding result： |
Contributor
Author
|
Hi @Tcc0403, could you please help me review my code? |
3 tasks
Tcc0403
requested changes
Jan 19, 2026
Collaborator
Tcc0403
left a comment
Tcc0403
left a comment
There was a problem hiding this comment.
It seems the current implementation is quite inefficient. I've left some comments about some possible issues it might have.
| ) | ||
|
|
||
|
|
||
| def get_optimal_block_size(total_elements, is_backward: bool): |
Collaborator
There was a problem hiding this comment.
what does is_backward do?
Comment on lines
+14
to
+24
| @triton.jit | ||
| def embedding_forward_kernel( | ||
| embeddings_ptr, | ||
| indices_ptr, | ||
| output_ptr, | ||
| total_elements, | ||
| n_elements, | ||
| embedding_dim: tl.constexpr, | ||
| BLOCK_SIZE: tl.constexpr, | ||
| NUM_STAGES: tl.constexpr, | ||
| ): |
Collaborator
There was a problem hiding this comment.
I think the original implementation with 2 block sizes for tile shape is more readable and more efficient.
persistant grid loop is fine, but the way this kernel loading embedding seems to be uncoalesced at some point.
Collaborator
There was a problem hiding this comment.
For instance, there will be some dim_idx not consecutive if BLOCK_SIZE is not multiple of embedding_dim. It will make the second tl.load trying to access different rows within a warp, as well as the last store.
Make these offsets created with 2d block size is more readable and efficient since we can avoid the uncoalesced access mentioned above.
Comment on lines
+110
to
+112
| tile_shapes = compute_default_tiling_strategy( | ||
| safety_margin=0.9, dtype_size=4, memory_multiplier=multiplier, shapes=((total_elements,),), tiling_dims=(0,) | ||
| ) |
Collaborator
There was a problem hiding this comment.
dtype_size should be embedding.dtype?
| block_size = tile_shapes[0][0] | ||
| return block_size | ||
| else: | ||
| return triton.next_power_of_2(total_elements) |
Collaborator
There was a problem hiding this comment.
I think fallback value should be workable, triton.next_power_of_2(total_elements) is too large.
| embeddings_ptr + embedding_offsets, | ||
| mask=final_mask, | ||
| other=0.0, | ||
| ).to(tl.float32) |
Collaborator
There was a problem hiding this comment.
any consideration why we need to upcast it?
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Summary
Add NPU support for the embedding.
Testing Done
I tested swiglu by following method and all cases passed:
python benchmark/scripts/benchmark_embedding.pypytest -v test/transformers/test_embedding.pymake testto ensure correctnessmake checkstyleto ensure code stylemake test-convergenceto ensure convergence