Skip to content

Commit 93d4f3b

Browse files
author
zhangyue
committed
refactor(pr66): trim commit-narration comments
/simplify found 4 comment blocks that narrate the rename rationale rather than encode load-bearing contracts: - `kernel_custom.h` forward-decl — compress build-system detail (`no_workspace_kernel`, `ascendc_library()`) to one line, keep only the ABI contract (`aclrtlaunch_<Entry>` is generated by AscendC from `op_kernel/`). - `op_host/<op>.cpp` `EXEC_KERNEL_CMD` — drop "Parameter order follows the base class: inputs, attributes, outputs."; the signature itself is self-evident. - `op_kernel/<op>.cpp` kernel entry — drop "Parameters follow the C2 convention ..." and "`aclrtlaunch_AddRmsNorm` matches the base `AddRmsNorm` class name"; these are commit-message material, not comments.
1 parent d60c180 commit 93d4f3b

6 files changed

Lines changed: 4 additions & 22 deletions

File tree

src/ascend/add_rms_norm/kernel_custom.h

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -14,11 +14,8 @@
1414
#include "base/add_rms_norm.h"
1515
#include "operator.h"
1616

17-
// Forward-declare the generated AscendC kernel launch function. This
18-
// symbol is provided by the `no_workspace_kernel` static library built
19-
// from `ascend/custom/add_rms_norm/op_kernel/add_rms_norm.cpp` via
20-
// `ascendc_library()`; the `aclrtlaunch_` prefix is prepended by the
21-
// AscendC toolchain to the kernel entry's `extern "C"` name.
17+
// Forward-declare the `aclrtlaunch_AddRmsNorm` launch symbol defined
18+
// by the AscendC toolchain from `custom/add_rms_norm/op_kernel/`.
2219
extern "C" uint32_t aclrtlaunch_AddRmsNorm(
2320
uint32_t block_dim, void* stream, void* input, void* residual, void* weight,
2421
int64_t total_rows, int64_t dim_length, int64_t dim_length_align,

src/ascend/custom/add_rms_norm/op_host/add_rms_norm.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -108,7 +108,6 @@ std::vector<at::Tensor> AddRmsNorm(const at::Tensor& x1, const at::Tensor& x2,
108108
// The first arg `AddRmsNorm` is the AscendC kernel entry-point name — it
109109
// must match the `__global__ __aicore__ void AddRmsNorm(...)` definition
110110
// in `op_kernel/` and the generated `aclrtlaunch_AddRmsNorm.h` header.
111-
// Parameter order follows the base class: inputs, attributes, outputs.
112111
EXEC_KERNEL_CMD(AddRmsNorm, block_dim, kernel_input1, kernel_input2,
113112
weight_float, total_rows, dim_length, dim_length_align,
114113
former_num, former_length, tail_length, eps_float,

src/ascend/custom/add_rms_norm/op_kernel/add_rms_norm.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -247,11 +247,6 @@ class KernelAddRmsNorm {
247247
// by the host launcher. fp16 and bf16 both have `sizeof == 2` but need
248248
// distinct numeric paths, so dispatch is on the `DataType` tag rather
249249
// than the byte size.
250-
//
251-
// Parameters follow the C2 convention: inputs first, attributes between,
252-
// outputs last. The kernel symbol is prefixed with `aclrtlaunch_` by the
253-
// `AscendC` toolchain, yielding `aclrtlaunch_AddRmsNorm` which matches the
254-
// base `AddRmsNorm` class name.
255250
extern "C" __global__ __aicore__ void AddRmsNorm(
256251
GM_ADDR input, GM_ADDR residual, GM_ADDR weight, int64_t total_rows,
257252
int64_t dim_length, int64_t dim_length_align, int64_t former_num,

src/ascend/custom/rms_norm/op_host/rms_norm.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,6 @@ at::Tensor RmsNorm(const at::Tensor& input, const at::Tensor& weight,
9797
// The first arg `RmsNorm` is the AscendC kernel entry-point name — it
9898
// must match the `__global__ __aicore__ void RmsNorm(...)` definition in
9999
// `op_kernel/` and the generated `aclrtlaunch_RmsNorm.h` header.
100-
// Parameter order follows the base class: inputs, attributes, outputs.
101100
EXEC_KERNEL_CMD(RmsNorm, block_dim, kernel_input, weight_float, total_rows,
102101
dim_length, dim_length_align, former_num, former_length,
103102
tail_length, eps_float, dtype_size_val, kernel_output);

src/ascend/custom/rms_norm/op_kernel/rms_norm.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -205,11 +205,6 @@ class KernelRmsNorm {
205205
// by the host launcher. fp16 and bf16 both have `sizeof == 2` but need
206206
// distinct numeric paths, so dispatch is on the `DataType` tag rather
207207
// than the byte size.
208-
//
209-
// Parameters follow the C2 convention: inputs first, attributes between,
210-
// outputs last. The kernel symbol is prefixed with `aclrtlaunch_` by the
211-
// `AscendC` toolchain, yielding `aclrtlaunch_RmsNorm` which matches the
212-
// base `RmsNorm` class name.
213208
extern "C" __global__ __aicore__ void RmsNorm(
214209
GM_ADDR input, GM_ADDR weight, int64_t total_rows, int64_t dim_length,
215210
int64_t dim_length_align, int64_t former_num, int64_t former_length,

src/ascend/rms_norm/kernel_custom.h

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -15,11 +15,8 @@
1515
#include "base/rms_norm.h"
1616
#include "operator.h"
1717

18-
// Forward-declare the generated AscendC kernel launch function. This
19-
// symbol is provided by the `no_workspace_kernel` static library built
20-
// from `ascend/custom/rms_norm/op_kernel/rms_norm.cpp` via
21-
// `ascendc_library()`; the `aclrtlaunch_` prefix is prepended by the
22-
// AscendC toolchain to the kernel entry's `extern "C"` name.
18+
// Forward-declare the `aclrtlaunch_RmsNorm` launch symbol defined by
19+
// the AscendC toolchain from `custom/rms_norm/op_kernel/`.
2320
extern "C" uint32_t aclrtlaunch_RmsNorm(
2421
uint32_t block_dim, void* stream, void* input, void* weight,
2522
int64_t total_rows, int64_t dim_length, int64_t dim_length_align,

0 commit comments

Comments
 (0)