fix: fix CUDA elementwise dynamic block shared memory

[chen2021673]

概要

修复 CUDA elementwise backward 在动态选择 block size 后，共享内存大小仍使用旧的模板参数 BLOCK_SIZE 计算的问题。

问题背景

LaunchKernel 已经改为根据 tensor size 通过 ChooseBlockSize() 动态选择 block size。但部分 elementwise backward launch 逻辑仍然使用模板参数 BLOCK_SIZE 来计算 dynamic shared memory。

这会导致实际 launch 配置和共享内存分配发生分歧：大 tensor 下运行时 block size 可能选择 512,但 shared memory 仍按 BLOCK_SIZE == 256 计算。bf16/half broadcast backward 的 block-reduction kernel 内部按 blockDim.x 访问 shared memory，最终 shared memory 分配不足，出现越界写。

修复方案

• 移除 elementwise launch helper 中不再可靠的 BLOCK_SIZE 模板参数，新增 ChooseBlockDims()，统一根据 ChooseBlockSize() 生成实际 launch block。
• forward/backward launch path 统一使用运行时选择的 block dims。
• dynamic shared memory 大小改为使用实际传入 launch lambda 的 block.x 计算。
• 新增 bf16 broadcast backward 大 block 回归测试。

回归测试

新增测试：

AutogradElementwiseBackwardTest.BFloat16MulBroadcastBackwardLargeBlock

覆盖场景：

• a: [512, 8192], bf16
• b: [8192], bf16 broadcast
• grad: [512, 8192], bf16

.62机器测试结果：

[kilinchange]

另外 commit message 麻烦修正下：
fix: fix CUDA elementwise dynamic block shared memory

[kilinchange]

1. 补充 node24 测试截图；
2. format 修复。

[chen2021673]

1. 补充 node24 测试截图；
2. format 修复。

1. node24 测试截图：

2. 已修复。