Skip to content

[hipblaslt] Update custom hhs/bbs tn kernel names #491

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
b-shi merged 2 commits into from
Jul 7, 2025
Merged

[hipblaslt] Update custom hhs/bbs tn kernel names #491

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
dc8b4e8
Update custom hhs/bbs tn kernel names
b-shi Jul 2, 2025
3f9f7f4
Merge branch 'develop' into custom_16b_org
b-shi Jul 7, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions ...e/Logic/asm_full/gfx950/Origami/gfx950_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18090,15 +18090,15 @@
AssertSummationElementMultiple: 1
AssignedDerivedParameters: true
AssignedProblemIndependentDerivedParameters: true
BaseName: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
BaseName: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
BufferLoad: true
BufferStore: true
CUCount: null
CUOccupancy: -1
ClusterLocalRead: 1
CodeObjectVersion: '4'
ConvertAfterDS: false
CustomKernelName: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
CustomKernelName: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
DebugStreamK: 0
DepthU: 64
DirectToLds: true
Expand Down Expand Up @@ -18133,7 +18133,7 @@
SupportUserGSU: false, UseUniversalArgs: true}
Kernel: true
KernelLanguage: Assembly
KernelNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
KernelNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
LDSTrInst: false
LSCA: 64
LSCB: 64
Expand Down Expand Up @@ -18236,7 +18236,7 @@
ScheduleIterAlg: 3
ScheduleLocalWrite: 1
SolutionIndex: 80
SolutionNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
SolutionNameMin: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
SourceSwap: false
StaggerU: 0
StaggerUMapping: 0
Expand Down
8 changes: 4 additions & 4 deletions ...e/Logic/asm_full/gfx950/Origami/gfx950_Cijk_Alik_Bljk_HHS_BH_BiasSH_HAS_SAV_UserArgs.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18090,15 +18090,15 @@
AssertSummationElementMultiple: 1
AssignedDerivedParameters: true
AssignedProblemIndependentDerivedParameters: true
BaseName: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
BaseName: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
BufferLoad: true
BufferStore: true
CUCount: null
CUOccupancy: -1
ClusterLocalRead: 1
CodeObjectVersion: '4'
ConvertAfterDS: false
CustomKernelName: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
CustomKernelName: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
DebugStreamK: 0
DepthU: 64
DirectToLds: true
Expand Down Expand Up @@ -18133,7 +18133,7 @@
SupportUserGSU: false, UseUniversalArgs: true}
Kernel: true
KernelLanguage: Assembly
KernelNameMin: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
KernelNameMin: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
LDSTrInst: false
LSCA: 64
LSCB: 64
Expand Down Expand Up @@ -18236,7 +18236,7 @@
ScheduleIterAlg: 3
ScheduleLocalWrite: 1
SolutionIndex: 80
SolutionNameMin: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
SolutionNameMin: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
SourceSwap: false
StaggerU: 0
StaggerUMapping: 0
Expand Down
14 changes: 7 additions & 7 deletions ...T256x256x64_MI16x16x1_shortname0_gfx950.s → ...T256x256x64_MI16x16x1_shortname0_gfx950.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,13 +4,13 @@
/******************************************/
.amdgcn_target "amdgcn-amd-amdhsa--gfx950"
.text
.protected Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.globl Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.protected Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.globl Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.p2align 8
.type Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950,@function
.type Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950,@function
.section .rodata,#alloc
.p2align 6
.amdhsa_kernel Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.amdhsa_kernel Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_accum_offset 256 // accvgpr offset
.amdhsa_next_free_vgpr 512 // vgprs
Expand Down Expand Up @@ -96,8 +96,8 @@ amdhsa.version:
- 1
- 1
amdhsa.kernels:
- .name: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.symbol: 'Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950.kd'
- .name: Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.symbol: 'Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950.kd'
.language: OpenCL C
.language_version:
- 2
Expand Down Expand Up @@ -338,7 +338,7 @@ amdhsa.kernels:
.wavefront_size: 64
...
.end_amdgpu_metadata
Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950:
Custom_Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950:
label_ASM_Start: /// Main body of the asm kernel
.macro V_MAGIC_DIV vgprDstIdx:req, dividend:req, magicNumber:req, magicShift:req, magicA:req
v_mul_hi_u32 v[\vgprDstIdx+1], \dividend, \magicNumber
Expand Down
14 changes: 7 additions & 7 deletions ...T256x256x64_MI16x16x1_shortname0_gfx950.s → ...T256x256x64_MI16x16x1_shortname0_gfx950.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,13 +4,13 @@
/******************************************/
.amdgcn_target "amdgcn-amd-amdhsa--gfx950"
.text
.protected Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.globl Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.protected Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.globl Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.p2align 8
.type Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950,@function
.type Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950,@function
.section .rodata,#alloc
.p2align 6
.amdhsa_kernel Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.amdhsa_kernel Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_accum_offset 256 // accvgpr offset
.amdhsa_next_free_vgpr 512 // vgprs
Expand Down Expand Up @@ -96,8 +96,8 @@ amdhsa.version:
- 1
- 1
amdhsa.kernels:
- .name: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.symbol: 'Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950.kd'
- .name: Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950
.symbol: 'Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950.kd'
.language: OpenCL C
.language_version:
- 2
Expand Down Expand Up @@ -338,7 +338,7 @@ amdhsa.kernels:
.wavefront_size: 64
...
.end_amdgpu_metadata
Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950:
Custom_Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_NTD_SK3_UserArgs_MT256x256x64_MI16x16x1_shortname0_gfx950:
label_ASM_Start: /// Main body of the asm kernel
.macro V_MAGIC_DIV vgprDstIdx:req, dividend:req, magicNumber:req, magicShift:req, magicA:req
v_mul_hi_u32 v[\vgprDstIdx+1], \dividend, \magicNumber
Expand Down