From 71ba123217ecf7ac24c665b10411e3b382e9f852 Mon Sep 17 00:00:00 2001 From: Koji Nakajima Date: Mon, 9 Mar 2026 04:55:06 +0000 Subject: [PATCH] [hipblaslt] Fix fail with kringshift.yaml --- .../tensilelite/Tensile/KernelWriter.py | 32 +++++++++---------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 3cb91083980..40b7b7edc9f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -2445,6 +2445,22 @@ def setupNewTile(self, kernel, tensorParametersA, tensorParametersB, isOptNLL=Fa module.addComment1("global read addresses: shift mxsb") module.add(self.graShift(kernel, tensorParametersB["MX"])) + # addresses + if not forceNoTileCode: + module.addComment1("global read addresses: addresses a") + module.add(self.graAddresses(kernel, tensorParametersA)) + if kernel["ProblemType"]["MXBlockA"]: + module.addComment1("global read addresses: addresses mxsa") + module.add(self.graAddresses(kernel, tensorParametersA["MX"])) + if kernel["ProblemType"]["Sparse"] and not kernel["DirectToVgprSparseMetadata"]: + module.addComment1("global read addresses: addresses metadata") + module.add(self.graAddresses(kernel, tPM)) + module.addComment1("global read addresses: addresses b") + module.add(self.graAddresses(kernel, tensorParametersB)) + if kernel["ProblemType"]["MXBlockB"]: + module.addComment1("global read addresses: addresses mxsb") + module.add(self.graAddresses(kernel, tensorParametersB["MX"])) + # workgoup SGPRs no longer needed module.add(self.removeGROffsetsVariableSgprsFromPool(kernel)) @@ -2468,22 +2484,6 @@ def setupNewTile(self, kernel, tensorParametersA, tensorParametersB, isOptNLL=Fa self.dontAppendCode = False self.dontAppendCode = self.dontAppendCode or forceNoTileCode - # addresses - if not forceNoTileCode: - module.addComment1("global read addresses: addresses a") - module.add(self.graAddresses(kernel, tensorParametersA)) - if kernel["ProblemType"]["MXBlockA"]: - module.addComment1("global read addresses: addresses mxsa") - module.add(self.graAddresses(kernel, tensorParametersA["MX"])) - if kernel["ProblemType"]["Sparse"] and not kernel["DirectToVgprSparseMetadata"]: - module.addComment1("global read addresses: addresses metadata") - module.add(self.graAddresses(kernel, tPM)) - module.addComment1("global read addresses: addresses b") - module.add(self.graAddresses(kernel, tensorParametersB)) - if kernel["ProblemType"]["MXBlockB"]: - module.addComment1("global read addresses: addresses mxsb") - module.add(self.graAddresses(kernel, tensorParametersB["MX"])) - # Add increment code gsuComponent = Component.GSU.find(self) module.add(gsuComponent.setupNewTile(self, kernel, tensorParametersA, tensorParametersB, tPM))