[NPU]: adjust MAX_FUSED_SIZE for NPU devices in group_norm (#1003)

- Set MAX_FUSED_SIZE to 16384 for NPU devices, keep 65536 for others
- Improve group_norm test cases with fixed, representative parameters
- Remove random test parameters in favor of deterministic test cases

Hardware Type: Ascend910B4
- [x] run `make test` to ensure correctness
- [x] run `make checkstyle` to ensure code style
- [ ] run `make test-convergence` to ensure convergence

Author: noemotiovon

src/liger_kernel/ops/group_norm.py

@@ -6,6 +6,7 @@
 
 from liger_kernel.ops.utils import compare_version
 from liger_kernel.ops.utils import ensure_contiguous
+from liger_kernel.utils import infer_device
 from liger_kernel.utils import is_npu_available
 
 if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
@@ -18,7 +19,10 @@
 else:
     from triton.language.math import rsqrt
 
-MAX_FUSED_SIZE = 65536
+if infer_device() == "npu":
+    MAX_FUSED_SIZE = 16384  # 8192
+else:
+    MAX_FUSED_SIZE = 65536
 
 
 @triton.jit

test/transformers/test_group_norm.py

@@ -1,5 +1,3 @@
-import random
-
 import pytest
 import torch
 
@@ -8,19 +6,15 @@
 
 device = infer_device()
 
-random_batch_size = random.randint(1, 16)
-random_num_groups = random.randint(1, 32)
-random_num_channels = random_num_groups * random.randint(1, 16)
-random_hidden_size = random.randint(1, 8192)
-
 
 @pytest.mark.parametrize(
     "batch_size, num_channels, num_groups, hidden_size",
     [
-        (1, 1, 1, 3),
-        (1, 4, 2, 4),
-        (16, 12, 3, 4096),
-        (random_batch_size, random_num_channels, random_num_groups, random_hidden_size),
+        (1, 1, 1, 3),  # minimal
+        (1, 32, 32, 4),  # group == channel
+        (16, 32, 1, 4096),  # single group
+        (2, 63, 21, 2163),  # non-aligned hidden
+        (16, 48, 12, 8192),  # large hidden
     ],
 )
 @pytest.mark.parametrize(