[Conv3D] Add packed-dot depthwise Conv3D hip kernel

[sammysun0711]

Integrates optimized FP16/BF16 packed-dot depthwise Conv3D kernel on gfx950, and FP16 packed-dot support on gfx942.

• For MI350X (gfx950), support both BF16 __builtin_amdgcn_fdot2_f32_bf16 and FP16 packed-dot __builtin_amdgcn_fdot2 LLVM built-in, Reduce VGPR from 86 to 47, use 68 SGPR, Reduce LDS usage from 32KB to 24KB, increase wave occupancy from 5 to 6. the packed dot BF16 kernel reaches +41.6% (BF16) and +40.9% (FP16) performance uplift compared to SGB while maintaining PyTorch correctness.
• For MI308X (gfx942), only support FP16 packed-dot __builtin_amdgcn_fdot2dot10-insts, #pragma unroll compiler codegen uses 255 VGPR, while limited LDS (64KB) does not improve wave occupancy (still 2), packed dot ISA improve perf by 34.9% while maintaining PyTorch correctness. Fallback to SGB for BF16 dtype since gfx942 does not support __builtin_amdgcn_fdot2_f32_bf16, which required dot12-insts.
• Selects it for supported gfx950/gfx942 workloads while preserving the existing SGB and original kernels as explicit
• Adds backend and dtype test controls, clearer benchmark labels, and wheel packaging for all contributed HIP sources.

MI350X (gfx950)

python test_conv_depthwise.py --depthwise-hip packed --dtype bf16

方法	平均耗时 (ms)	吞吐量 (TFLOPS)	带宽 (GB/s)	预热/Tune (ms)	Max Diff
Standard PyTorch Conv3d	6.0155	2.71	73.55	61.11	-
run_conv_depthwise_3d_pyhip (original)	0.6465	25.23	684.36	7.76	0.000004
run_conv_depthwise_3d_pyhip (sgb)	0.5745	28.39	770.09	6.75	0.000004
run_conv_depthwise_3d_pyhip (packed)	0.4059	40.19	1090.12	4.76	0.000004

python test_conv_depthwise.py --depthwise-hip packed --dtype fp16

方法	平均耗时 (ms)	吞吐量 (TFLOPS)	带宽 (GB/s)	预热/Tune (ms)	Max Diff
Standard PyTorch Conv3d	5.9082	2.76	74.89	61.38	-
run_conv_depthwise_3d_pyhip (original)	0.6622	24.63	668.12	8.03	0.000000
run_conv_depthwise_3d_pyhip (sgb)	0.5824	28.01	759.72	7.05	0.000000
run_conv_depthwise_3d_pyhip (packed)	0.4133	39.47	1070.59	4.83	0.000000

MI308X (gfx942)

python test_conv_depthwise.py --depthwise-hip packed --dtype fp16

方法	平均耗时 (ms)	吞吐量 (TFLOPS)	带宽 (GB/s)	预热/Tune (ms)	Max Diff
Standard PyTorch Conv3d	28.5808	0.57	15.48	285.86	-
run_conv_depthwise_3d_pyhip (original)	3.1337	5.21	141.19	31.66	0.000000
run_conv_depthwise_3d_pyhip (sgb)	3.1965	5.10	138.42	32.29	0.000000
run_conv_depthwise_3d_pyhip (packed)	2.3717	6.88	186.55	23.95	0.000000

[Copilot]

Pull request overview

This PR integrates a new packed-dot depthwise Conv3D HIP kernel and adds an automatic kernel-selection path so PyHIP can choose the best implementation for supported AMD GPU architectures (gfx950/gfx942) and dtypes, alongside benchmark/test and packaging updates.

Changes:

• Added a new packed-dot depthwise Conv3D HIP kernel implementation (conv_depthwise3d_hip_packed_dot.cpp).
• Implemented hip_impl="auto" selection logic in conv_depthwise.py, choosing packed when supported and otherwise falling back to sgb.
• Extended the benchmark/test harness with dtype/backend controls and improved output labeling; updated packaging to include contributed .cpp sources in wheels.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 3 comments.

File	Description
tests/contrib/conv3d/test_conv_depthwise.py	Adds --dtype and --depthwise-hip auto/packed/sgb/original controls and adjusts benchmark labels/formatting.
src/contrib/conv_depthwise3d_hip_packed_dot.cpp	New packed-dot depthwise Conv3D kernel using fdot2 builtins and reduced LDS/VGPR strategy.
src/contrib/conv_depthwise.py	Adds packed/auto backend resolution based on shape, dtype, and detected AMDGPU arch; switches default to hip_impl="auto".
pyproject.toml	Ensures pyhip.contrib ships .cpp sources as package data for wheel builds.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.