fix(tensilelite): revert Stream-K launch limit changes (#12022)

Revert the Stream-K launch limit changes from #11518.

TheRock Multi-Arch CI showed that the rocm-libraries #8183 bump failed the
standard TensileLite test job on both gfx94X and gfx950. Restoring the three
files changed by #11518 on top of that failing pin made both targeted jobs pass.

Fixes #12021

Author: jayhawk-commits

projects/hipblaslt/tensilelite/src/ContractionSolution.cpp

@@ -5454,9 +5454,7 @@ namespace TensileLite
                 {
                     // Limit workgroups per CU to 3
                     // TODO Verify this limit is best
-                    const size_t occupancy
-                        = std::max(self.sizeMapping.CUOccupancy, static_cast<int>(1));
-                    auto kernelOccupancy = std::min(occupancy, size_t{3});
+                    auto kernelOccupancy = std::min(self.sizeMapping.CUOccupancy, 3);
                     auto maxGrid         = cuCount * kernelOccupancy;
                     if(pAMDGPU->skMaxCUs > 0)
                     {
@@ -5680,22 +5678,11 @@ namespace TensileLite
             {
                 size_t itersPerTile = problem.getItersPerTile(self.sizeMapping);
                 size_t itersPerWG   = tiles * itersPerTile / skGrid;
-                // numWorkItems = skGrid * threadsPerWorkGroup
-                // the largest grid this fallback can express is
-                // UINT32_MAX / threadsPerWorkGroup -- 2^24 tiles at 256 threads
-                // but 2^25 at 128 and 2^22 at 1024.
-                const size_t threadsPerWorkGroup = self.sizeMapping.workGroupSize.x
-                                                    * self.sizeMapping.workGroupSize.y
-                                                    * self.sizeMapping.workGroupSize.z;
-                const size_t maxTiles = (size_t{std::numeric_limits<uint32_t>::max()} + 1)
-                                        / std::max(threadsPerWorkGroup, size_t{1});
 
                 if(itersPerTile >= 65536 || itersPerWG >= 65536
-                   || (tiles * itersPerTile) >= maxTiles)
+                   || (tiles * itersPerTile) >= 16777216)
                 {
-                    const size_t occupancy
-                    = std::max(self.sizeMapping.CUOccupancy, static_cast<int>(1));
-                    skGrid = (tiles >= maxTiles) ? cuCount * occupancy : tiles;
+                    skGrid = tiles;
                     if(outTreeBoundsFallback)
                         *outTreeBoundsFallback = true;
                 }