issue/826: kunlun layernorm #827
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
xgqdut2016
reviewed
Dec 23, 2025
| __global__ void layerNormKernel( | ||
| int32_t loop_idx, | ||
| Tdata *output, // [b, seq, dim] | ||
| Tdata *output_standardization, // [b, seq, dim] |
Collaborator
There was a problem hiding this comment.
这几个注释[b, seq, dim]是不是删掉比较好，给人一种只能处理3D向量的错觉
| int32_t offset_output_standardization = 0; | ||
| int32_t offset_output_rstd_deviation = 0; | ||
| int32_t offset_input = 0; | ||
| for (int i = 0; i < ndim - 1; i++) { |
Collaborator
There was a problem hiding this comment.
这个循环和上面计算t_coords[]的循环。类似能否合并
| int32_t offset_output_rstd_deviation = 0; | ||
| int32_t offset_input = 0; | ||
| for (int i = 0; i < ndim - 1; i++) { | ||
| int32_t dim_i = shape_local[i].value; |
Collaborator
There was a problem hiding this comment.
这个循环和上面计算t_coords[]的循环能不能合并成一个循环，比如说也不要单独申请内存存储t_coords，直接使用寄存器变量，借助一次循环得到input和output的index，这样代码以及速度会不会更好一些，还是说有什么特殊考虑
| ? 255 | ||
| : static_cast<int32_t>(info.othersize); | ||
| int32_t num_loops = (static_cast<int32_t>(info.othersize) + num_blocks - 1) / num_blocks; | ||
| for (int32_t i = 0; i < num_loops; i++) { |
Collaborator
There was a problem hiding this comment.
能不能把这个任务分配过程挪到kernel里面，这样写循环，会导致增加很多kernel启动开销吧
xgqdut2016
approved these changes
Dec 23, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Uh oh!
There was an error while loading. Please reload this page.