rocmPackages: 7.1.1 -> 7.2.0 (#484792)

This commit is contained in:
Luna Nova
2026-03-05 13:38:58 +00:00
committed by GitHub
54 changed files with 442 additions and 399 deletions
@@ -21,12 +21,12 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "amdsmi";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "rocm";
repo = "amdsmi";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-BGe3+8YFwu41ZVAF+VtN5Cn9pfzGxmCg/Rpq8qWOEoM=";
hash = "sha256-3V3B9+B3cpg0ebvsmX6wwvKoTIHbXtvbpIQHs4tkeWU=";
};
postPatch = ''
@@ -8,12 +8,12 @@
stdenv.mkDerivation {
pname = "aqlprofile";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "aqlprofile";
tag = "rocm-7.1.1";
tag = "rocm-7.2.0";
hash = "sha256-MAZUHo52gb0aZSFnKugMlXxcDkmMwhy1AFF1RDRgRVk=";
};
@@ -70,7 +70,7 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "clr";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -84,7 +84,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "clr";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-ofsq1uqMixtum5C6cp/UgTDpgGPfj+rAd6PoDx5iLLw=";
hash = "sha256-zz2O4Qsl1zXMC25L714azsFR2PROAvdpjgKhRolmt1w=";
};
nativeBuildInputs = [
@@ -148,8 +148,8 @@ stdenv.mkDerivation (finalAttrs: {
./cmake-find-x11-libgl.patch
(fetchpatch {
# [PATCH] rocclr: Extend HIP ISA compatibility checks
sha256 = "sha256-InUSIFI1MgkfocBEoZjO2BCgXNyfF10ehh9jkTtAPXs=";
url = "https://github.com/GZGavinZhao/rocm-systems/commit/937dcfdd316b589509c061809186fe5451d22431.patch";
hash = "sha256-3MsDL+OQg24wH1RDhbao74RuIbzEAmduwla9KOPzQ/M=";
url = "https://github.com/GZGavinZhao/rocm-systems/commit/039cb23b24d739adb8c0f9de8b550d9f557de031.patch";
relative = "projects/clr";
})
];
@@ -42,7 +42,7 @@ stdenv.mkDerivation (finalAttrs: {
'';
pname = "composable_kernel_base";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -58,7 +58,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "composable_kernel";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-exdkyTIK03dzlCXHm3j8ehEb9NxLOxPX9QyfMsiCgSs=";
hash = "sha256-ABL0MSmWtqAeY5uyw8Ib64npB2v82baUnzLpmrEgDn4=";
};
nativeBuildInputs = [
@@ -54,6 +54,7 @@ let
targets = [
"device_grouped_conv1d_fwd_instance"
"device_grouped_conv2d_fwd_instance"
"device_grouped_conv2d_fwd_bias_bnorm_clamp_instance"
"device_grouped_conv2d_fwd_bias_clamp_instance"
"device_grouped_conv2d_fwd_clamp_instance"
"device_grouped_conv2d_fwd_dynamic_op_instance"
@@ -77,6 +78,7 @@ let
targets = [
"device_grouped_conv3d_fwd_instance"
"device_grouped_conv3d_fwd_clamp_instance"
"device_grouped_conv3d_fwd_bias_bnorm_clamp_instance"
"device_grouped_conv3d_fwd_bias_clamp_instance"
"device_grouped_conv3d_fwd_bilinear_instance"
"device_grouped_conv3d_fwd_convinvscale_instance"
@@ -9,7 +9,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "half";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
@@ -7,13 +7,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "hip-common";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "HIP";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-aI41HdyzyE5HtCbzFmTQYGMiFyQELM/MTVOAbQsg8So=";
hash = "sha256-eQ+jHc6MlZePPIwJQMB8NKiEcE26i83+U21vKgfGFFM=";
};
dontConfigure = true;
@@ -8,13 +8,13 @@
}:
stdenv.mkDerivation (finalAttrs: {
pname = "hipblas-common";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "hipBLAS-common";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-sswqYIv3c5ArT6dQ3MRru7ZZVA0VxoKW/a0pYTBpjgE=";
hash = "sha256-oOelDrk7TLe9/17fzaw6CDPL4MGAoITJL90ahHzhaAE=";
};
nativeBuildInputs = [
@@ -24,7 +24,7 @@
# Can also use cuBLAS
stdenv.mkDerivation (finalAttrs: {
pname = "hipblas";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -43,7 +43,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipBLAS";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-2wDnYZapJ4fU/JZpy6qOf5+DYuatWmCw/hA8WO5x2T0=";
hash = "sha256-w7myy4FO2vKm91OOfSGWO3bxR8jAvjU68oqI7aJX82c=";
};
postPatch = ''
@@ -1,53 +1,3 @@
diff --git a/tensilelite/Tensile/SolutionStructs/Naming.py b/tensilelite/Tensile/SolutionStructs/Naming.py
index 4f220960db1d..99535e246650 100644
--- a/tensilelite/Tensile/SolutionStructs/Naming.py
+++ b/tensilelite/Tensile/SolutionStructs/Naming.py
@@ -105,7 +105,6 @@ def _getName(state, requiredParameters: frozenset, splitGSU: bool, ignoreInterna
if splitGSU:
state["GlobalSplitU"] = "M" if (state["GlobalSplitU"] > 1 or state["GlobalSplitU"] == -1) else state["GlobalSplitU"]
-
requiredParametersTemp = set(requiredParameters.union(["GlobalSplitU"]))
if ignoreInternalArgs:
diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py
index bab8c687509..e03f456fbec 100644
--- a/tensilelite/Tensile/CustomYamlLoader.py
+++ b/tensilelite/Tensile/CustomYamlLoader.py
@@ -1,3 +1,6 @@
+# Copyright © Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
import yaml
from pathlib import Path
Author: Luna Nova <git@lunnova.dev>
Date: Sun Oct 12 11:52:10 2025 -0700
[hipblaslt] intern strings to reduce duplicate memory for solution keys
diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py
index 685e69220c..9fdf38d8e5 100644
--- a/tensilelite/Tensile/CustomYamlLoader.py
+++ b/tensilelite/Tensile/CustomYamlLoader.py
@@ -1,6 +1,7 @@
# Copyright © Advanced Micro Devices, Inc., or its affiliates.
# SPDX-License-Identifier: MIT
+import sys
import yaml
from pathlib import Path
@@ -85,7 +86,7 @@ def parse_scalar(loader: yaml.Loader):
if not evt.style:
return None
- return value
+ return sys.intern(value)
def load_yaml_stream(yaml_path: Path, loader_type: yaml.Loader):
with open(yaml_path, 'r') as f:
diff --git a/tensilelite/Tensile/Common/Parallel.py b/tensilelite/Tensile/Common/Parallel.py
index 1a2bf9e119..f46100c7b8 100644
--- a/tensilelite/Tensile/Common/Parallel.py
@@ -55,7 +5,7 @@ index 1a2bf9e119..f46100c7b8 100644
@@ -22,43 +22,58 @@
#
################################################################################
-import concurrent.futures
-import itertools
+import multiprocessing
@@ -67,10 +17,10 @@ index 1a2bf9e119..f46100c7b8 100644
-from joblib import Parallel, delayed
+from functools import partial
+from typing import Any, Callable
from .Utilities import tqdm
-def joblibParallelSupportsGenerator():
- import joblib
- from packaging.version import Version
@@ -81,7 +31,7 @@ index 1a2bf9e119..f46100c7b8 100644
+ return int(os.environ['CMAKE_BUILD_PARALLEL_LEVEL'])
+ except ValueError:
+ pass
- joblibVer = joblib.__version__
- return Version(joblibVer) >= Version("1.4.0")
+ # 2. Parse MAKEFLAGS for -jN
@@ -89,12 +39,12 @@ index 1a2bf9e119..f46100c7b8 100644
+ match = re.search(r'-j\s*(\d+)', makeflags)
+ if match:
+ return int(match.group(1))
+ return -1
-def CPUThreadCount(enable=True):
- from .GlobalParameters import globalParameters
+def CPUThreadCount(enable=True):
if not enable:
return 1
@@ -121,7 +71,7 @@ index 1a2bf9e119..f46100c7b8 100644
- if cpuThreads == -1:
- return cpu_count
+ cpuThreads = len(os.sched_getaffinity(0))
- return min(cpu_count, cpuThreads)
+ if os.name == "nt":
+ # Windows supports at most 61 workers because the scheduler uses
@@ -129,13 +79,13 @@ index 1a2bf9e119..f46100c7b8 100644
+ # is actually 64, but some handles are needed for accounting).
+ cpuThreads = min(cpuThreads, 61)
+ return max(1, cpuThreads)
def pcallWithGlobalParamsMultiArg(f, args, newGlobalParameters):
@@ -71,19 +86,22 @@ def pcallWithGlobalParamsSingleArg(f, arg, newGlobalParameters):
return f(arg)
-def apply_print_exception(item, *args):
- # print(item, args)
+def OverwriteGlobalParameters(newGlobalParameters):
@@ -166,8 +116,8 @@ index 1a2bf9e119..f46100c7b8 100644
finally:
@@ -98,154 +116,121 @@ def OverwriteGlobalParameters(newGlobalParameters):
GlobalParameters.globalParameters.update(newGlobalParameters)
-def ProcessingPool(enable=True, maxTasksPerChild=None):
- import multiprocessing
- import multiprocessing.dummy
@@ -191,10 +141,10 @@ index 1a2bf9e119..f46100c7b8 100644
+def progress_logger(iterable, total, message, min_log_interval=5.0):
+ """
+ Generator that wraps an iterable and logs progress with time-based throttling.
+ Only logs progress if at least min_log_interval seconds have passed since last log.
+ Only prints completion message if task took >= min_log_interval seconds.
-def ParallelMap(function, objects, message="", enable=True, method=None, maxTasksPerChild=None):
+ Yields (index, item) tuples.
"""
@@ -212,7 +162,7 @@ index 1a2bf9e119..f46100c7b8 100644
+ start_time = time.time()
+ last_log_time = start_time
+ log_interval = 1 + (total // 100)
- threadCount = CPUThreadCount(enable)
- pool = ProcessingPool(enable, maxTasksPerChild)
-
@@ -251,12 +201,12 @@ index 1a2bf9e119..f46100c7b8 100644
+ print(f"{message}\t{idx+1: 5d}/{total: 5d}")
+ last_log_time = current_time
+ yield idx, item
- if message != "":
- message += ": "
+ elapsed = time.time() - start_time
+ final_idx = idx + 1 if 'idx' in locals() else 0
- print("{0}Launching {1} threads{2}...".format(message, threadCount, countMessage))
- sys.stdout.flush()
- currentTime = time.time()
@@ -268,8 +218,8 @@ index 1a2bf9e119..f46100c7b8 100644
- return rv
+ if elapsed >= min_log_interval or last_log_time > start_time:
+ print(f"{message} done in {elapsed:.1f}s!\t{final_idx: 5d}/{total: 5d}")
-def ParallelMapReturnAsGenerator(function, objects, message="", enable=True, multiArg=True):
- from .GlobalParameters import globalParameters
+def imap_with_progress(pool, func, iterable, total, message, chunksize):
@@ -277,10 +227,10 @@ index 1a2bf9e119..f46100c7b8 100644
+ for _, result in progress_logger(pool.imap(func, iterable, chunksize=chunksize), total, message):
+ results.append(result)
+ return results
- threadCount = CPUThreadCount(enable)
- print("{0}Launching {1} threads...".format(message, threadCount))
- if threadCount <= 1 and globalParameters["ShowProgressBar"]:
- # Provide a progress bar for single-threaded operation.
- callFunc = lambda args: function(*args) if multiArg else lambda args: function(args)
@@ -288,7 +238,7 @@ index 1a2bf9e119..f46100c7b8 100644
+def _ParallelMap_generator(worker, objects, objLen, message, chunksize, threadCount, globalParameters, maxtasksperchild):
+ # separate fn because yield makes the entire fn a generator even if unreachable
+ ctx = multiprocessing.get_context('forkserver' if os.name != 'nt' else 'spawn')
- with concurrent.futures.ProcessPoolExecutor(max_workers=threadCount) as executor:
- resultFutures = (executor.submit(function, *arg if multiArg else arg) for arg in objects)
- for result in concurrent.futures.as_completed(resultFutures):
@@ -297,8 +247,8 @@ index 1a2bf9e119..f46100c7b8 100644
+ initializer=OverwriteGlobalParameters, initargs=(globalParameters,)) as pool:
+ for _, result in progress_logger(pool.imap_unordered(worker, objects, chunksize=chunksize), objLen, message):
+ yield result
def ParallelMap2(
- function, objects, message="", enable=True, multiArg=True, return_as="list", procs=None
+ function: Callable,
@@ -333,7 +283,7 @@ index 1a2bf9e119..f46100c7b8 100644
"""
- Generally equivalent to list(map(function, objects)), possibly executing in parallel.
+ from .GlobalParameters import globalParameters
- message: A message describing the operation to be performed.
- enable: May be set to false to disable parallelism.
- multiArg: True if objects represent multiple arguments
@@ -342,23 +292,23 @@ index 1a2bf9e119..f46100c7b8 100644
- if return_as in ("generator", "generator_unordered") and not joblibParallelSupportsGenerator():
- return ParallelMapReturnAsGenerator(function, objects, message, enable, multiArg)
+ threadCount = CPUThreadCount(enable)
- from .GlobalParameters import globalParameters
+ if not hasattr(objects, "__len__"):
+ objects = list(objects)
- threadCount = procs if procs else CPUThreadCount(enable)
+ objLen = len(objects)
+ if objLen == 0:
+ return [] if return_as == "list" else iter([])
- threadCount = CPUThreadCount(enable)
+ f = (lambda x: function(*x)) if multiArg else function
+ if objLen == 1:
+ print(f"{message}: (1 task)")
+ result = [f(x) for x in objects]
+ return result if return_as == "list" else iter(result)
- if threadCount <= 1 and globalParameters["ShowProgressBar"]:
- # Provide a progress bar for single-threaded operation.
- return [function(*args) if multiArg else function(args) for args in tqdm(objects, message)]
@@ -367,7 +317,7 @@ index 1a2bf9e119..f46100c7b8 100644
+ if objLen
+ else ""
+ )
- countMessage = ""
- try:
- countMessage = " for {} tasks".format(len(objects))
@@ -419,19 +369,19 @@ index ffceb636f5..127b3386a1 100644
--- a/tensilelite/Tensile/CustomKernels.py
+++ b/tensilelite/Tensile/CustomKernels.py
@@ -24,7 +24,9 @@
from . import CUSTOM_KERNEL_PATH
from Tensile.Common.ValidParameters import checkParametersAreValid, validParameters, newMIValidParameters
+from Tensile.CustomYamlLoader import DEFAULT_YAML_LOADER
+from functools import lru_cache
import yaml
import os
@@ -58,10 +60,13 @@ def getCustomKernelConfigAndAssembly(name, directory=CUSTOM_KERNEL_PATH):
return (config, assembly)
+# getCustomKernelConfig will get called repeatedly on the same file
+# 20x logic loading speedup for aquavanjaram_Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HAS_SAB_SAV_freesize_custom_GSUs
+@lru_cache
@@ -442,13 +392,34 @@ index ffceb636f5..127b3386a1 100644
+ return yaml.load(rawConfig, Loader=DEFAULT_YAML_LOADER)["custom.config"]
except yaml.scanner.ScannerError as e:
raise RuntimeError("Failed to read configuration for custom kernel: {0}\nDetails:\n{1}".format(name, e))
diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py
index e03f456fbe..ed7510f2ce 100644
--- a/tensilelite/Tensile/CustomYamlLoader.py
+++ b/tensilelite/Tensile/CustomYamlLoader.py
@@ -1,6 +1,7 @@
# Copyright © Advanced Micro Devices, Inc., or its affiliates.
# SPDX-License-Identifier: MIT
+import sys
import yaml
from pathlib import Path
@@ -70,7 +71,7 @@ def parse_scalar(loader: yaml.Loader):
elif is_float(value_lower):
return float(value_lower)
- return value
+ return sys.intern(value)
def load_yaml_stream(yaml_path: Path, loader_type: yaml.Loader):
with open(yaml_path, 'r') as f:
diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py
index 835ed9c019..024c6c49c1 100644
index 22d19851a3..348068b3bf 100644
--- a/tensilelite/Tensile/TensileCreateLibrary/Run.py
+++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py
@@ -26,8 +26,10 @@ import rocisa
import functools
import glob
+import gc
@@ -460,8 +431,8 @@ index 835ed9c019..024c6c49c1 100644
from timeit import default_timer as timer
@@ -78,6 +80,25 @@ from Tensile.Utilities.Decorators.Timing import timing
from .ParseArguments import parseArguments
+def getMemoryUsage():
+ """Get peak and current memory usage in MB."""
+ rusage = resource.getrusage(resource.RUSAGE_SELF)
@@ -484,11 +455,11 @@ index 835ed9c019..024c6c49c1 100644
class KernelCodeGenResult(NamedTuple):
err: int
src: str
@@ -115,6 +136,29 @@ def processKernelSource(kernelWriterAssembly, data, splitGSU, kernel) -> KernelC
@@ -115,6 +136,29 @@ def processKernelSource(kernelWriterAssembly, data, outOptions, splitGSU, kernel
)
+def processAndAssembleKernelTCL(kernelWriterAssembly, rocisa_data, splitGSU, kernel, assemblyTmpPath, assembler):
+def processAndAssembleKernelTCL(kernelWriterAssembly, rocisa_data, outOptions, splitGSU, kernel, assemblyTmpPath, assembler):
+ """
+ Pipeline function for TCL mode that:
+ 1. Generates kernel source
@@ -496,7 +467,7 @@ index 835ed9c019..024c6c49c1 100644
+ 3. Assembles to .o file
+ 4. Deletes .s file
+ """
+ result = processKernelSource(kernelWriterAssembly, rocisa_data, splitGSU, kernel)
+ result = processKernelSource(kernelWriterAssembly, rocisa_data, outOptions, splitGSU, kernel)
+ return writeAndAssembleKernel(result, assemblyTmpPath, assembler)
+
+
@@ -516,8 +487,8 @@ index 835ed9c019..024c6c49c1 100644
removeKernelNames = []
@@ -189,6 +233,24 @@ def writeAssembly(asmPath: Union[Path, str], result: KernelCodeGenResult):
return path, isa, wfsize, minResult
+def writeAndAssembleKernel(result: KernelCodeGenResult, asmPath: Union[Path, str], assembler):
+ """Write assembly file and immediately assemble it to .o file"""
+ if result.err:
@@ -539,13 +510,14 @@ index 835ed9c019..024c6c49c1 100644
def writeHelpers(
outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H
):
@@ -268,13 +330,14 @@ def writeSolutionsAndKernels(
@@ -272,14 +334,15 @@ def writeSolutionsAndKernels(
numAsmKernels = len(asmKernels)
numKernels = len(asmKernels)
assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite"
- asmIter = zip(
- itertools.repeat(kernelWriterAssembly),
- itertools.repeat(rocisa.rocIsa.getInstance().getData()),
- itertools.repeat(outOptions),
- itertools.repeat(splitGSU),
- asmKernels
+
@@ -553,6 +525,7 @@ index 835ed9c019..024c6c49c1 100644
+ processKernelSource,
+ kernelWriterAssembly=kernelWriterAssembly,
+ data=rocisa.rocIsa.getInstance().getData(),
+ outOptions=outOptions,
+ splitGSU=splitGSU
)
- asmResults = ParallelMap2(processKernelSource, asmIter, "Generating assembly kernels", return_as="list")
@@ -560,12 +533,12 @@ index 835ed9c019..024c6c49c1 100644
removeInvalidSolutionsAndKernels(
asmResults, asmKernels, solutions, errorTolerant, getVerbosity(), splitGSU
)
@@ -282,19 +345,21 @@ def writeSolutionsAndKernels(
@@ -287,19 +350,21 @@ def writeSolutionsAndKernels(
asmResults, asmKernels, solutions, splitGSU
)
- def assemble(ret):
- p, isa, wavefrontsize, result = ret
- p, isa, wavefrontsize, _ = ret
- asmToolchain.assembler(isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o")))
-
- unaryWriteAssembly = functools.partial(writeAssembly, assemblyTmpPath)
@@ -586,33 +559,41 @@ index 835ed9c019..024c6c49c1 100644
)
+ del asmResults
+ gc.collect()
writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H)
srcKernelFile = Path(outputPath) / "Kernels.cpp"
@@ -369,32 +434,31 @@ def writeSolutionsAndKernelsTCL(
@@ -376,40 +441,35 @@ def writeSolutionsAndKernelsTCL(
uniqueAsmKernels = [k for k in asmKernels if not k.duplicate]
- def assemble(ret):
- p, isa, wavefrontsize, result = ret
- asmToolchain.assembler(isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o")))
- def assemble(ret, removeTemporaries: bool):
- asmPath, isa, wavefrontsize, result = ret
- asmToolchain.assembler(isaToGfx(isa), wavefrontsize, str(asmPath), str(asmPath.with_suffix(".o")))
- if removeTemporaries:
- asmPath.unlink()
- return result
-
- unaryAssemble = functools.partial(assemble, removeTemporaries=removeTemporaries)
-
outOptions = rocisa.rocIsa.getInstance().getOutputOptions()
outOptions.outputNoComment = not disableAsmComments
- unaryProcessKernelSource = functools.partial(
- processKernelSource,
+ processKernelFn = functools.partial(
+ processAndAssembleKernelTCL,
kernelWriterAssembly,
rocisa.rocIsa.getInstance().getData(),
outOptions,
splitGSU,
+ assemblyTmpPath=assemblyTmpPath,
+ assembler=asmToolchain.assembler
)
- unaryWriteAssembly = functools.partial(writeAssembly, assemblyTmpPath)
- compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F)
- ret = ParallelMap2(
- compose(assemble, unaryWriteAssembly, unaryProcessKernelSource),
- compose(unaryAssemble, unaryWriteAssembly, unaryProcessKernelSource),
+ results = ParallelMap2(
+ processKernelFn,
uniqueAsmKernels,
@@ -635,10 +616,10 @@ index 835ed9c019..024c6c49c1 100644
buildAssemblyCodeObjectFiles(
asmToolchain.linker,
asmToolchain.bundler,
@@ -493,6 +557,15 @@ def generateKernelHelperObjects(solutions: List[Solution], cxxCompiler: str, isa
@@ -508,6 +568,15 @@ def generateKernelHelperObjects(solutions: List[Solution], cxxCompiler: str, isa
return sorted(khos, key=sortByEnum, reverse=True) # Ensure that we write Enum kernel helpers are first in list
+def libraryIter(lib: MasterSolutionLibrary):
+ if len(lib.solutions):
+ for i, s in enumerate(lib.solutions.items()):
@@ -650,11 +631,11 @@ index 835ed9c019..024c6c49c1 100644
+
@timing
def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInfoMap):
@@ -508,26 +581,23 @@ def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInf
@@ -523,26 +592,23 @@ def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInf
printSolutionRejectionReason = True
printIndexAssignmentInfo = False
- fIter = zip(
- logicFiles,
- itertools.repeat(assembler),
@@ -672,7 +653,7 @@ index 835ed9c019..024c6c49c1 100644
+ isaInfoMap=isaInfoMap,
+ lazyLibraryLoading=args["LazyLibraryLoading"]
)
- def libraryIter(lib: MasterSolutionLibrary):
- if len(lib.solutions):
- for i, s in enumerate(lib.solutions.items()):
@@ -691,31 +672,31 @@ index 835ed9c019..024c6c49c1 100644
+ multiArg=False,
):
_, architectureName, _, _, _, newLibrary = library
@@ -539,6 +609,9 @@ def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInf
@@ -554,6 +620,9 @@ def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInf
else:
masterLibraries[architectureName] = newLibrary
masterLibraries[architectureName].version = args["CodeObjectVersion"]
+ del library, newLibrary
+
+ gc.collect()
# Sort masterLibraries to make global soln index values deterministic
solnReIndex = 0
@@ -734,6 +807,9 @@ def run():
@@ -751,6 +820,9 @@ def run():
)
stop_wsk = timer()
print(f"Time to generate kernels (s): {(stop_wsk-start_wsk):3.2f}")
+ numKernelHelperObjs = len(kernelHelperObjs)
+ del kernelWriterAssembly, kernelHelperObjs
+ gc.collect()
archs = [ # is this really different than the other archs above?
isaToGfx(arch)
@@ -751,13 +827,10 @@ def run():
@@ -768,13 +840,10 @@ def run():
if kName not in solDict:
solDict["%s"%kName] = kernel
- def writeMsl(name, lib):
- filename = os.path.join(newLibraryDir, name)
- lib.applyNaming(splitGSU)
@@ -725,13 +706,13 @@ index 835ed9c019..024c6c49c1 100644
LibraryIO.write(filename, libraryMapping, "msgpack")
+ del libraryMapping
+ gc.collect()
start_msl = timer()
for archName, newMasterLibrary in masterLibraries.items():
@@ -774,12 +847,22 @@ def run():
@@ -791,12 +860,22 @@ def run():
kName = getKeyNoInternalArgs(s.originalSolution, splitGSU)
s.sizeMapping.CUOccupancy = solDict["%s"%kName]["CUOccupancy"]
- ParallelMap2(writeMsl,
+ writeFn = functools.partial(
+ writeMasterSolutionLibrary,
@@ -749,15 +730,15 @@ index 835ed9c019..024c6c49c1 100644
print(f"Time to write master solution libraries (s): {(stop_msl-start_msl):3.2f}")
+ del masterLibraries, solutions, kernels, solDict
+ gc.collect()
if not arguments["KeepBuildTmp"]:
buildTmp = Path(arguments["OutputPath"]).parent / "library" / "build_tmp"
@@ -796,8 +879,11 @@ def run():
@@ -813,8 +892,11 @@ def run():
print("")
stop = timer()
+ peak_memory_mb, current_memory_mb = getMemoryUsage()
print(f"Total time (s): {(stop-start):3.2f}")
print(f"Total kernels processed: {numKernels}")
print(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}")
@@ -772,7 +753,7 @@ index e33c617b6f..ba163e9918 100644
@@ -303,8 +303,7 @@ def avoidRegressions(originalDir, incrementalDir, outputPath, forceMerge, noEff=
logicsFiles[origFile] = origFile
logicsFiles[incFile] = incFile
- iters = zip(logicsFiles.keys())
- logicsList = ParallelMap2(loadData, iters, "Loading Logics...", return_as="list")
+ logicsList = ParallelMap2(loadData, logicsFiles.keys(), "Loading Logics...", return_as="list", multiArg=False)
@@ -785,11 +766,11 @@ index 5ff265d0ed..c1803a6349 100644
+++ b/tensilelite/Tensile/TensileUpdateLibrary.py
@@ -26,7 +26,7 @@ from . import LibraryIO
from .Tensile import addCommonArguments, argUpdatedGlobalParameters
from .Common import assignGlobalParameters, print1, restoreDefaultGlobalParameters, HR, \
- globalParameters, architectureMap, ensurePath, ParallelMap, __version__
+ globalParameters, architectureMap, ensurePath, ParallelMap2, __version__
import argparse
import copy
@@ -149,7 +149,7 @@ def TensileUpdateLibrary(userArgs):
@@ -798,8 +779,8 @@ index 5ff265d0ed..c1803a6349 100644
fIter = zip(logicFiles, itertools.repeat(args.logic_path), itertools.repeat(outputPath))
- libraries = ParallelMap(UpdateLogic, fIter, "Updating logic files", method=lambda x: x.starmap)
+ libraries = ParallelMap2(UpdateLogic, fIter, "Updating logic files", multiArg=True, return_as="list")
def main():
diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py
index a8b91e8d62..265e1d532c 100644
@@ -808,16 +789,16 @@ index a8b91e8d62..265e1d532c 100644
@@ -30,7 +30,7 @@ import subprocess
from pathlib import Path
from typing import List, Union, NamedTuple
-from Tensile.Common import print2
+from Tensile.Common import print1, print2
from Tensile.Common.Architectures import isaToGfx
from ..SolutionStructs import Solution
@@ -92,8 +92,26 @@ def buildAssemblyCodeObjectFiles(
if coName:
coFileMap[asmDir / (coName + extCoRaw)].add(str(asmDir / (kernel["BaseName"] + extObj)))
+ # Build reference count map for .o files to handle shared object files
+ # (.o files from kernels marked .duplicate in TensileCreateLibrary)
+ objFileRefCount = collections.Counter()
@@ -70,14 +70,17 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "hipblaslt${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocm-libraries";
rev = "a676499add42941ff6af1e8d3f0504416dac7429";
hash = "sha256-zIYdHFbHyP2V6dkx6Ueb6NBqWu8tJji2hSWF9zWEJa4=";
sparseCheckout = [ "projects/hipblaslt" ];
rev = "rocm-${finalAttrs.version}";
hash = "sha256-I2dGn4Ld5lZeML8GePcLPssplBZ+4weNR6uBEqFdZVg=";
sparseCheckout = [
"projects/hipblaslt"
"shared"
];
};
sourceRoot = "${finalAttrs.src.name}/projects/hipblaslt";
env.CXX = compiler;
@@ -203,7 +206,7 @@ stdenv.mkDerivation (finalAttrs: {
# Move binaries to appropriate outputs and delete leftover /bin
+ ''
mkdir -p $benchmark/bin
mv $out/bin/hipblaslt-{api-overhead,sequence,bench*} $out/bin/*.yaml $out/bin/*.py $benchmark/bin
mv $out/bin/hipblaslt-{api-overhead,bench*} $out/bin/*.yaml $out/bin/*.py $benchmark/bin
${lib.optionalString buildTests ''
mkdir -p $test/bin
mv $out/bin/hipblas-test $test/bin
@@ -17,7 +17,7 @@
# CUB can also be used as a backend instead of rocPRIM.
stdenv.mkDerivation (finalAttrs: {
pname = "hipcub";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -33,7 +33,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipCUB";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-3/499pX7hFKJwbhSEH3/bMhNdyabK36jLv8s5TC2fuw=";
hash = "sha256-hIIpU1E/X6bPIhFfauNLPnFNBYSwtoxZV335Huw9fCA=";
};
nativeBuildInputs = [
@@ -21,7 +21,7 @@
# Can also use cuFFT
stdenv.mkDerivation (finalAttrs: {
pname = "hipfft";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -40,7 +40,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipFFT";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-6FyI9s6H/lHFhm8aUqB9+vvJ0CRIVWCFLNoJprvsI6o=";
hash = "sha256-5zeL3o/Dfh0iyZOj5CHyfEQdFNehD2BLmllnZRdpsQI=";
fetchSubmodules = true;
};
@@ -1,7 +1,6 @@
{
lib,
stdenv,
fetchpatch,
fetchFromGitHub,
rocmUpdateScript,
cmake,
@@ -11,24 +10,15 @@
stdenv.mkDerivation (finalAttrs: {
pname = "hipfort";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "hipfort";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-IVLhp8rYtKhkt8K8Mc0qyrp670oKoK0QeclnJjO36pY=";
hash = "sha256-XaB4jauCN41tgD1YHHA2td/yckwfMBemBe/iL0SCxQo=";
};
patches = [
(fetchpatch {
name = "hipfort-fix-cmake-4.patch";
url = "https://github.com/ROCm/hipfort/commit/75552c7ec48e3bd6a914c57c9475ec573ccb37d9.patch";
hash = "sha256-S9r1V6cUo9QbKbu/NK4wIvXMV6BFv7+/n9mjCScVk40=";
includes = [ "bin/*" ];
})
];
nativeBuildInputs = [
cmake
rocm-cmake
@@ -12,13 +12,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "hipify";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "HIPIFY";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-NbFFHAAvMGpvIryhEbktN5w03Cpay9lEqelqkUT9dpQ=";
hash = "sha256-LC0lnYetV7RPVw92zew6za6bDH4zmnERXUM4MVaRVtc=";
};
strictDeps = true;
@@ -14,7 +14,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "hiprand";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -27,7 +27,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipRAND";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-yypKwJ5p2aSAqapEHre9eOjjaFAI8Q1dDYWhwfrrdhg=";
hash = "sha256-mhhwYswewcRYKGoVAB/ZTIY8ubFSlTDXKZFoJpoDf2o=";
};
nativeBuildInputs = [
@@ -21,7 +21,7 @@
# Can also use cuSOLVER
stdenv.mkDerivation (finalAttrs: {
pname = "hipsolver";
version = "7.1.0";
version = "7.2.0";
outputs = [
"out"
@@ -40,7 +40,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipSOLVER";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-n0+Xy92mBhI2vjB7pFHGt2UgPIHo9rO2ZxGxkz30Otw=";
hash = "sha256-xrzRU3X0zlmK6t8Bz4p5ihnJnadrqiW1BW2HLRDpu1M=";
};
nativeBuildInputs = [
@@ -19,7 +19,7 @@
# This can also use cuSPARSE as a backend instead of rocSPARSE
stdenv.mkDerivation (finalAttrs: {
pname = "hipsparse";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -35,7 +35,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "hipSPARSE";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-zaMAsLfUYkMUiA2vnGiK0//+vZwugqt5T6xJU4FBrRo=";
hash = "sha256-DmanHHlDny/SweYrSes4xQeWoF6TV+AoThyDHQpz+hQ=";
};
nativeBuildInputs = [
+60 -30
View File
@@ -2,7 +2,7 @@
lib,
stdenv,
# LLVM version closest to ROCm fork to override
llvmPackages_20,
llvmPackages_22,
overrideCC,
lndir,
rocm-device-libs,
@@ -22,23 +22,27 @@
# leaving compressed line tables (-g1 -gz) unstripped
# TODO: Should also apply to downstream packages which use rocmClangStdenv?
profilableStdenv ? false,
# FIXME: proper two-stage bootstrap & PGO/BOLT/LTO
# LTO currently disabled due to llvm 22 vs 22.1 bitcode mismatch between
# llvmPackages_22 (22.1) and ROCm LLVM (22.0). Uses new bitcode attr
# 105 nocreateundeforpoison fails to link hipify.
# Whether to use LTO when building the ROCm toolchain
# Slows down this toolchain's build, for typical ROCm usecase
# time saved building composable_kernel and other heavy packages
# will outweight that. ~3-4% speedup multiplied by thousands
# of corehours.
withLto ? true,
withLto ? false,
# whether rocm stdenv uses libcxx (clang c++ stdlib) instead of gcc stdlibc++
withLibcxx ? false,
}:
let
version = "7.1.1";
version = "7.2.0";
# major version of this should be the clang version ROCm forked from
rocmLlvmVersion = "20.0.0-rocm";
rocmLlvmVersion = "22.0.0-rocm";
# llvmPackages_base version should match rocmLlvmVersion
# so libllvm's bitcode is compatible with the built toolchain
llvmPackages_base = llvmPackages_20;
llvmPackages_base = llvmPackages_22;
llvmPackagesNoBintools = llvmPackages_base.override {
bootBintools = null;
bootBintoolsNoLibc = null;
@@ -115,7 +119,7 @@ let
owner = "ROCm";
repo = "llvm-project";
rev = "rocm-${version}";
hash = "sha256-CfknIRVeR1bCKh1xzXKl3ehVp0kWT0uGrI9C1HTSKVo=";
hash = "sha256-I/Bxq1JjU9N4h3vXj/tbD3xPYY4+N7QzYA8UTIq1EL0=";
};
llvmMajorVersion = lib.versions.major rocmLlvmVersion;
# An llvmPackages (pkgs/development/compilers/llvm/) built from ROCm LLVM's source tree
@@ -197,10 +201,23 @@ let
rm -rf $out/lib/cmake $out/lib/lib*.a
mkdir -p $out/lib/clang/${llvmMajorVersion}/lib/linux/
ln -s $out/lib/linux/libclang_rt.* $out/lib/clang/${llvmMajorVersion}/lib/linux/
mkdir -p $out/lib/clang/${llvmMajorVersion}/lib/${stdenv.hostPlatform.config}/
for f in $out/lib/linux/*-${stdenv.hostPlatform.parsed.cpu.name}.*; do
[ -e "$f" ] || continue
base="$(basename "$f" | sed 's/-${stdenv.hostPlatform.parsed.cpu.name}\././')"
ln -sf "$f" "$out/lib/clang/${llvmMajorVersion}/lib/${stdenv.hostPlatform.config}/$base"
done
find $out -type f -exec sed -i "s|${cc.out}|$out|g" {} +
find $out -type f -exec sed -i "s|${cc.dev}|$out|g" {} +
# Clang config file: redirect resource dir to the sysroot (where compiler-rt
# lives) and set GCC install prefix for header/library search
cat > $out/bin/${stdenv.hostPlatform.config}.cfg <<CFG
-resource-dir=$out/lib/clang/${llvmMajorVersion}
${lib.optionalString (!withLibcxx) "--gcc-toolchain=${gcc-prefix}"}
CFG
${lib.getExe rdfind} -makesymlinks true ${
builtins.concatStringsSep " " (map (x: "${x}/lib") paths)
} $out/ # create links *within* the sysroot to save space
@@ -233,6 +250,9 @@ let
(lib.cmakeFeature "LLVM_ENABLE_ZSTD" "FORCE_ON")
# required for threaded ThinLTO to work
(lib.cmakeBool "LLVM_ENABLE_THREADS" true)
# third-party/benchmark is broken in rocm-7.2.0
# error: '__COUNTER__' is a C2y extension [-Werror,-Wc2y-extensions]
(lib.cmakeBool "LLVM_INCLUDE_BENCHMARKS" false)
# LLVM tries to call git to embed VCS info if FORCE_VC_ aren't set
(lib.cmakeFeature "LLVM_FORCE_VC_REVISION" "rocm-${version}")
(lib.cmakeFeature "LLVM_FORCE_VC_REPOSITORY" "https://github.com/ROCm/llvm-project")
@@ -276,11 +296,14 @@ overrideLlvmPackagesRocm (s: {
libllvm = (s.prev.libllvm.override { }).overrideAttrs (old: {
patches = old.patches ++ [
(fetchpatch {
# fix compile error in tools/gold/gold-plugin.cpp
name = "gold-plugin-fix.patch";
url = "https://github.com/llvm/llvm-project/commit/b0baa1d8bd68a2ce2f7c5f2b62333e410e9122a1.patch";
hash = "sha256-yly93PvGIXOnFeDGZ2W+W6SyhdWFM6iwA+qOeaptrh0=";
relative = "llvm";
# Revert of a patch that cause perf regression by pessimising loop unrolling decisions
# See PR and issue discussion
# https://github.com/ROCm/rocm-systems/issues/2865 https://github.com/ROCm/llvm-project/pull/1349
name = "rocm-llvm-revert-unrolling-regression.patch";
url = "https://github.com/ROCm/llvm-project/commit/f58b06dce1f9c15707c5f808fd002e18c2accf7e.patch";
hash = "sha256-pH+3C7PSDqNfOF014sA5Rvm+sc2IJMQJfysS2bvj/o0=";
# stripLen instead of relative to avoid filterdiff mangling /dev/null on the deleted test file
stripLen = 1;
})
./perf-increase-namestring-size.patch
# TODO: consider reapplying "Don't include aliases in RegisterClassInfo::IgnoreCSRForAllocOrder"
@@ -342,7 +365,17 @@ overrideLlvmPackagesRocm (s: {
passthru = old.passthru // {
inherit gcc-prefix;
};
patches = old.patches ++ [
patches = [
(fetchpatch {
# [clang][cmake] Add option to control scan-build-py installation (#172727)
name = "clang-scan-build-py-configurable.patch";
url = "https://github.com/llvm/llvm-project/commit/f5759eeb63a3a5ce7d555c13c3126cea84e0c7b1.patch";
relative = "clang";
hash = "sha256-73IDPGZWKX4vny3x5FJ3/NQw8XRad9UNwfYkvQdMB4s=";
})
]
++ old.patches
++ [
# Never add FHS include paths
./clang-bodge-ignore-systemwide-incls.diff
# Prevents builds timing out if a single compiler invocation is very slow but
@@ -352,11 +385,23 @@ overrideLlvmPackagesRocm (s: {
./perf-shorten-gcclib-include-paths.patch
(fetchpatch {
# [ClangOffloadBundler]: Add GetBundleIDsInFile to OffloadBundler
sha256 = "sha256-G/mzUdFfrJ2bLJgo4+mBcR6Ox7xGhWu5X+XxT4kH2c8=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/6d296f879b0fed830c54b2a9d26240da86c8bb3a.patch";
hash = "sha256-OsarDZXuJ5vAXTP4i0NBUeK/r6tQPumaqmMWkf29UtM=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/c7de294b0d1d25f277f9d1cbb2c9e09c7600e210.patch";
relative = "clang";
})
];
# ROCm 7.2 commits 4dda51261a6 "Replace hostexec with upstream rpc"
# and 2ca1509d6d2 "Put the RTL, Back!", added CGEmitEmissaryExec.cpp which
# includes ../../openmp/device/include/EmissaryIds.h, breaking
# standalone clang builds. The upstream PR llvm/llvm-project#175265
# ("[OpenMP] support for Emissary APIs") moves EmissaryIds.h into
# clang/lib/Headers/ so this should not be needed once that lands
# and ROCm rebases onto it.
postUnpack = ''
${old.postUnpack or ""}
mkdir -p "''${sourceRoot}/openmp/device/include"
ln -s "${llvmSrc}/openmp/device/include/EmissaryIds.h" "''${sourceRoot}/openmp/device/include/"
'';
hardeningDisable = [ "all" ];
nativeBuildInputs = old.nativeBuildInputs ++ [
removeReferencesTo
@@ -375,14 +420,7 @@ overrideLlvmPackagesRocm (s: {
# https://github.com/llvm/llvm-project/blob/6976deebafa8e7de993ce159aa6b82c0e7089313/clang/cmake/caches/DistributionExample-stage2.cmake#L9-L11
cmakeFlags =
# TODO: Remove in followup, tblgen now works correctly but would rebuild
(builtins.filter tablegenUsage old.cmakeFlags)
++ commonCmakeFlags
++ lib.optionals (!withLibcxx) [
# FIXME: Config file in rocm-toolchain instead of GCC_INSTALL_PREFIX?
# Expected to be fully removed eventually
"-DUSE_DEPRECATED_GCC_INSTALL_PREFIX=ON"
"-DGCC_INSTALL_PREFIX=${gcc-prefix}"
];
(builtins.filter tablegenUsage old.cmakeFlags) ++ commonCmakeFlags;
preFixup = ''
${toString old.preFixup or ""}
moveToOutput "lib/lib*.a" "$dev"
@@ -437,14 +475,6 @@ overrideLlvmPackagesRocm (s: {
isGNU = false;
};
compiler-rt-libc = s.prev.compiler-rt-libc.overrideAttrs (old: {
patches = old.patches ++ [
# fix build with glibc >= 2.42
(fetchpatch {
url = "https://github.com/llvm/llvm-project/commit/59978b21ad9c65276ee8e14f26759691b8a65763.patch";
hash = "sha256-ys5SMLfO3Ay9nCX9GV5yRCQ6pLsseFu/ZY6Xd6OL4p0=";
relative = "compiler-rt";
})
];
meta = old.meta // llvmMeta;
});
compiler-rt = s.final.compiler-rt-libc;
@@ -17,6 +17,6 @@ index 33f12fb2d075d..ba35f65210c7a 100644
if (TCArgs.hasFlag(options::OPT_offload_compress,
- options::OPT_no_offload_compress, false))
+ options::OPT_no_offload_compress, true))
CmdArgs.push_back("-compress");
CmdArgs.push_back("--compress");
if (TCArgs.hasArg(options::OPT_v))
CmdArgs.push_back("-verbose");
@@ -1,4 +1,4 @@
From ef6c5b353861be727c98f1319c81d0c6b609d644 Mon Sep 17 00:00:00 2001
From 8c64deff2a82fcca5b521362193a1fd8b1a0dca7 Mon Sep 17 00:00:00 2001
From: Luna Nova <git@lunnova.dev>
Date: Tue, 17 Dec 2024 04:29:11 -0800
Subject: [PATCH] HACK: Get canonical GCC include path so doesn't have
@@ -6,21 +6,20 @@ Subject: [PATCH] HACK: Get canonical GCC include path so doesn't have
This allows more of the strings used in compilation to fit inside
fixed size stack allocated buffers instead of spilling into the heap
---
clang/lib/Driver/ToolChains/Gnu.cpp | 16 ++++++++++------
lib/Driver/ToolChains/Gnu.cpp | 16 ++++++++++------
1 file changed, 10 insertions(+), 6 deletions(-)
diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp
index af9fd46f0ce7b..a63a7a93f6a78 100644
diff --git a/lib/Driver/ToolChains/Gnu.cpp b/lib/Driver/ToolChains/Gnu.cpp
index 07df9864636..b1201b999f0 100644
--- a/lib/Driver/ToolChains/Gnu.cpp
+++ b/lib/Driver/ToolChains/Gnu.cpp
@@ -3394,29 +3394,33 @@ bool Generic_GCC::addLibStdCXXIncludePaths(Twine IncludeDir, StringRef Triple,
if (!getVFS().exists(IncludeDir))
@@ -3326,29 +3326,33 @@ static bool addLibStdCXXIncludePaths(llvm::vfs::FileSystem &vfs,
if (!vfs.exists(IncludeDir))
return false;
+ SmallString<260> CanonicalIncludeDir;
+ if (getVFS().getRealPath(IncludeDir, CanonicalIncludeDir))
+ if (vfs.getRealPath(IncludeDir, CanonicalIncludeDir))
+ return false;
+
// Debian native gcc uses g++-multiarch-incdir.diff which uses
@@ -34,24 +33,27 @@ index af9fd46f0ce7b..a63a7a93f6a78 100644
- (Include + "/" + Triple + Dir.substr(Include.size()) + IncludeSuffix)
+ (Include + "/" + Triple + CanonicalIncludeDir.substr(Include.size()) + IncludeSuffix)
.str();
if (DetectDebian && !getVFS().exists(Path))
if (DetectDebian && !vfs.exists(Path))
return false;
// GPLUSPLUS_INCLUDE_DIR
- addSystemInclude(DriverArgs, CC1Args, IncludeDir);
+ addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir);
- ToolChain::addSystemInclude(DriverArgs, CC1Args, IncludeDir);
+ ToolChain::addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir);
// GPLUSPLUS_TOOL_INCLUDE_DIR. If Triple is not empty, add a target-dependent
// include directory.
if (DetectDebian)
addSystemInclude(DriverArgs, CC1Args, Path);
ToolChain::addSystemInclude(DriverArgs, CC1Args, Path);
else if (!Triple.empty())
addSystemInclude(DriverArgs, CC1Args,
- IncludeDir + "/" + Triple + IncludeSuffix);
+ CanonicalIncludeDir + "/" + Triple + IncludeSuffix);
ToolChain::addSystemInclude(DriverArgs, CC1Args,
- IncludeDir + "/" + Triple + IncludeSuffix);
+ CanonicalIncludeDir + "/" + Triple + IncludeSuffix);
// GPLUSPLUS_BACKWARD_INCLUDE_DIR
- addSystemInclude(DriverArgs, CC1Args, IncludeDir + "/backward");
+ if (getVFS().exists(CanonicalIncludeDir + "/backward"))
+ addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir + "/backward");
- ToolChain::addSystemInclude(DriverArgs, CC1Args, IncludeDir + "/backward");
+ if (vfs.exists(CanonicalIncludeDir + "/backward"))
+ ToolChain::addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir + "/backward");
return true;
}
--
2.52.0
@@ -58,7 +58,7 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "migraphx";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -74,7 +74,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "AMDMIGraphX";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-s6w4bF7koK4wnf6leVKBzwIB4X2ROHa3EgX6XuJIAew=";
hash = "sha256-FBAVsk4x3ATLDYtfzcPUPnUTwe36maIAY1/FwqG6jD0=";
};
nativeBuildInputs = [
@@ -4,7 +4,7 @@
callPackage,
fetchFromGitHub,
fetchpatch,
rocmUpdateScript,
fetchurl,
runCommand,
pkg-config,
cmake,
@@ -43,7 +43,7 @@
let
# FIXME: cmake files need patched to include this properly
cFlags = "-Wno-documentation-pedantic --offload-compress -I${hipblas-common}/include -I${hipblas}/include -I${roctracer}/include -I${nlohmann_json}/include -I${sqlite.dev}/include -I${rocrand}/include";
version = "7.1.1";
version = "7.2.0";
# Targets outside this list will get
# error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
@@ -65,42 +65,12 @@ let
"gfx1201"
] gpuTargets;
kdbTargets = lib.intersectLists [
"gfx900"
"gfx906"
"gfx908"
"gfx90a"
"gfx942"
"gfx1030"
] gpuTargets;
src = fetchFromGitHub {
owner = "ROCm";
repo = "MIOpen";
rev = "rocm-${version}";
hash = "sha256-g0AEpuUiwKCu/doiRh9mW34W04m4ynHoarSyl6tR/aE=";
fetchLFS = true;
hash = "sha256-OwBFzUruHHeJD7n3zLs/NsU5cNEjwkwFgim4m2/1hR0=";
fetchSubmodules = true;
# WORKAROUND: .lfsconfig is incorrectly set to exclude everything upstream
leaveDotGit = true;
# FIXME: if someone can reduce the level of awful here that would be really nice
postFetch = ''
export HOME=$(mktemp -d)
cd $out
git remote add origin $url
git fetch origin +refs/tags/rocm-${version}:refs/tags/rocm-${version}
git clean -fdx
git switch -c rocm-${version} refs/tags/rocm-${version}
git config lfs.fetchexclude "none"
rm .lfsconfig
git lfs install
git lfs track "*.kdb.bz2"
git lfs fetch --include="src/kernels/**"
git lfs pull --include="src/kernels/**"
git lfs checkout
rm -rf .git
'';
};
latex = lib.optionalAttrs buildDocs (
@@ -120,19 +90,27 @@ let
)
);
kernelDatabases = lib.genAttrs kdbTargets (
target:
runCommand "miopen-${target}.kdb" { preferLocalBuild = true; } ''
${lbzip2}/bin/lbzip2 -ckd ${src}/src/kernels/${target}.kdb.bz2 > $out
''
);
# Kernel databases moved from Git LFS to DVC (anonymous s3 bucket s3://therock-dvc/rocm-libraries)
fetchKdb =
name:
{ url, hash }:
runCommand "miopen-${name}.kdb" { preferLocalBuild = true; } ''
${lbzip2}/bin/lbzip2 -ckd ${
fetchurl {
inherit url hash;
name = "${name}.kdb.bz2";
}
} > $out
'';
kdbs = lib.mapAttrs fetchKdb (import ./kdbs.nix);
linkKDBsTo =
targetPath:
lib.concatStringsSep "" (
map (target: ''
ln -sf ${kernelDatabases.${target}} ${targetPath}/${target}.kdb
'') kdbTargets
lib.mapAttrsToList (name: kdb: ''
ln -sf ${kdb} ${targetPath}/${name}.kdb
'') kdbs
);
in
stdenv.mkDerivation (finalAttrs: {
@@ -153,14 +131,13 @@ stdenv.mkDerivation (finalAttrs: {
url = "https://github.com/ROCm/MIOpen/commit/3413d2daaeb44b7d6eadcc03033a5954a118491e.patch";
hash = "sha256-ST4snUcTmmSI1Ogx815KEX9GdMnmubsavDzXCGJkiKs=";
})
# FIXME: We need to rebase or drop this arch compat patch
# https://github.com/ROCm/MIOpen/issues/3540 suggests that
# arch compat patching doesn't work correctly for gfx1031
# (fetchpatch {
# name = "Extend-MIOpen-ISA-compatibility.patch";
# url = "https://github.com/GZGavinZhao/MIOpen/commit/416088b534618bd669a765afce59cfc7197064c1.patch";
# hash = "sha256-OwONCA68y8s2GqtQj+OtotXwUXQ5jM8tpeM92iaD4MU=";
# })
(fetchpatch {
# [miopen] Extend HIP ISA compatibility
name = "Extend-MIOpen-ISA-compatibility.patch";
url = "https://github.com/GZGavinZhao/rocm-libraries/commit/02f0fedffdc197f146dd45f41e10990a00cde3ee.patch";
hash = "sha256-My32iZw75rvB4fyvUJJ2kw2bU9/39awGteFGjzijixw=";
relative = "projects/miopen";
})
];
outputs = [
@@ -308,11 +285,7 @@ stdenv.mkDerivation (finalAttrs: {
inherit frugally-deep nlohmann_json;
};
};
passthru.updateScript = rocmUpdateScript {
name = finalAttrs.pname;
inherit (finalAttrs.src) owner;
inherit (finalAttrs.src) repo;
};
passthru.updateScript = ./update.sh;
meta = {
description = "Machine intelligence library for ROCm";
@@ -0,0 +1,27 @@
# Generated by update.sh
{
gfx1030 = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/22/c4bd38cc34d3906add061ee57b099b";
hash = "sha256-LESExp7sFl7osCM7HdVMpzOE0dU7953e2yzgeHSGgfI=";
};
gfx900 = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/cf/9c12dafe25b15ea1e93088b7b061af";
hash = "sha256-lFFB5GQFTjUnx3JDx816yH+96OGRt9XvBzfLFFHtUrE=";
};
gfx906 = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/04/b3be6a15c700ff382e0a5e136010e1";
hash = "sha256-cn1ODLcUQi+/QJUlWbELukMUj4+b1vjjRhhTseL9tdM=";
};
gfx908 = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/35/b9082b72e661dc5e37f14de1d9d4ed";
hash = "sha256-pI30zpZL3OEP3aeLdVpM00nHeTyD5K+XDH47g0W2A+M=";
};
gfx90a = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/a1/1615323e06c5774fe19cf34c51baef";
hash = "sha256-TkzSgY3V/Xs5YXfr69TrL8JjqDr1MZMgeBeQ/e2kEjE=";
};
gfx942 = {
url = "https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5/6a/fe97d7bd441fe15280b28b6e3b1748";
hash = "sha256-fyhnR3+xPFkHLN36dI9UYIhqoVstt5RpareGn2S9lMM=";
};
}
+48
View File
@@ -0,0 +1,48 @@
#!/usr/bin/env nix-shell
#!nix-shell -i bash -p curl jq nix common-updater-scripts
set -euo pipefail
gh_curl() { curl ${GITHUB_TOKEN:+-u ":$GITHUB_TOKEN"} -sf "$@"; }
version="${1:-$(gh_curl "https://api.github.com/repos/ROCm/MIOpen/releases?per_page=4" \
| jq -re 'map(.tag_name // .name)
| map(select(test("^rocm-[0-9]+\\.[0-9]+(\\.[0-9]+)?$")))
| first | ltrimstr("rocm-")')}"
# main source version
echo "Updating MIOpen to $version" >&2
update-source-version rocmPackages.miopen "$version" --ignore-same-hash
# kdbs.nix from s3 bucket
cd "$(dirname "${BASH_SOURCE[0]}")"
bucket="https://therock-dvc.s3.amazonaws.com/rocm-libraries/files/md5"
# Fetch directory listing; extract name + download_url for .dvc files
dvc_entries=$(gh_curl "https://api.github.com/repos/ROCm/MIOpen/contents/src/kernels?ref=rocm-${version}" \
| jq -r '.[] | select(.name | endswith(".kdb.bz2.dvc")) | "\(.name)\t\(.download_url)"')
[[ -n "$dvc_entries" ]] || { echo "No .kdb.bz2.dvc files found, upstream likely changed packaging again" >&2; exit 1; }
{
echo "# Generated by update.sh"
echo "{"
while IFS=$'\t' read -r dvc_file raw_url; do
arch="${dvc_file%.kdb.bz2.dvc}"
md5=$(curl -sf "$raw_url" | awk '/- md5:/{print $3}')
url="${bucket}/${md5:0:2}/${md5:2}"
sri=$(nix --extra-experimental-features nix-command store prefetch-file --json "$url" --name "${arch}.kdb.bz2" | jq -r .hash)
echo " $arch: $sri" >&2
cat <<EOF
$arch = {
url = "$url";
hash = "$sri";
};
EOF
done <<< "$dvc_entries"
echo "}"
} > kdbs.nix
echo "Wrote kdbs.nix" >&2
@@ -40,13 +40,13 @@ stdenv.mkDerivation (finalAttrs: {
"cpu"
);
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "MIVisionX";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-aC6GUPK6ZSOAx+PHHB4gLPNAG5U/kapqX7YWilusDw8=";
hash = "sha256-rdYQt001kFi9FjVL5R8XVzT+Y+rZVeMsbjXCJEB887A=";
};
patches = [
@@ -40,7 +40,7 @@ in
# infiniband ib_peer_mem support isn't in the mainline kernel but is carried by some distros
stdenv.mkDerivation (finalAttrs: {
pname = "rccl${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -58,7 +58,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rccl";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-3u7D3Gre1n+4Lf+cK+RMfCUM9c46pXZjdhGOrwIKM0w=";
hash = "sha256-Kpobn0fR7/mXGVWYjNQ35sToRWLfOQTHRbjuSYj1of4=";
};
requiredSystemFeatures = [ "big-parallel" ]; # Very resource intensive LTO
@@ -49,7 +49,7 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rdc";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -65,7 +65,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rdc";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-oJS0fBeISpgjZkaMfFl/Yq7mPOsdaZDqiRjSPv7kT1Q=";
hash = "sha256-gF17mcbsqsL2K9pwTAVG9E6nVY4pi2lC88WyPa0xMes=";
};
patches = [
@@ -22,7 +22,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocalution";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -41,7 +41,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocALUTION";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-UkXbcbhaTulCqW1FZHwTyUZLFSnj7WmLYqlPqu5m6YM=";
hash = "sha256-7ZlY0GVIsUJE8OOAkQwQ4pXNHKzgCQomQOET9+TKQN4=";
};
nativeBuildInputs = [
@@ -40,13 +40,13 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rocblas${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocBLAS";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-obKypbYmnSeOtOr7g0pxOz02YfzZ0bGugTtznkeHz14=";
hash = "sha256-oY6yELFnpnDDksu63Go88TGLAgC64tUaPIbaRgD2qus=";
};
outputs = [ "out" ] ++ lib.optional buildBenchmarks "benchmark" ++ lib.optional buildTests "test";
@@ -137,9 +137,9 @@ stdenv.mkDerivation (finalAttrs: {
patches = [
(fetchpatch {
name = "Extend-rocBLAS-HIP-ISA-compatibility.patch";
url = "https://github.com/GZGavinZhao/rocm-libraries/commit/2850f22f80f90c9e498f520608a82989a4932ec3.patch";
hash = "sha256-SPsdEwGe+r8bQudkChRzBDAgu3tPQWFweZCgzh+4nOE=";
stripLen = 2;
url = "https://github.com/GZGavinZhao/rocm-libraries/commit/49f21f3cfe7eb4b8ac724eef81fa2cae97a3c22e.patch";
hash = "sha256-vAVVpTwt49lGHu2YopR1X68v5LwFOjUYuSC4ucBpFGg=";
relative = "projects/rocblas";
})
];
@@ -47,7 +47,7 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rocdbgapi";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -62,7 +62,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "ROCdbgapi";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-RwYZJPwGhNtSSvmSgy0AsNTc98cav0/u9jH5f93sB9M=";
hash = "sha256-KqvhwfIv8pbr8WbnfAKl71fg5yxbwYcpzZcGU9Htdkc=";
};
nativeBuildInputs = [
@@ -20,13 +20,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocfft${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocFFT";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-1Ho3b5NmnzfLhDnvH6FECigs6OgpbJrxw4EnqKyaHA0=";
hash = "sha256-AFBV5tlhNYH9+d0KG7mCjSifKfRt+8c1ie5wASTiQXM=";
};
nativeBuildInputs = [
@@ -23,13 +23,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocgdb";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "ROCgdb";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-LDVGO++voqwVMM6RcRtgbXTUvFLTyg/TFdSZanv5Xdc=";
hash = "sha256-oml/HLExnnjh7+axeWPZRWecpwK2BnzVOaGvXYhrxKs=";
};
nativeBuildInputs = [
@@ -9,13 +9,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocm-cmake";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocm-cmake";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-im6UO0crO0Jc27zkTsdvJYPHit8IGlw/vDPGrmP1XqY=";
hash = "sha256-gY6jzIIN1pSXGbCMN6y35Q/VJgbIqWDRjD8aI/fc1L0=";
};
nativeBuildInputs = [ cmake ];
@@ -30,20 +30,20 @@ stdenv.mkDerivation (finalAttrs: {
patches = [
# [Comgr] Extend ISA compatibility
(fetchpatch {
sha256 = "sha256-dgow0kwSWM1TnkqWOZDRQrh5nuF8p5jbYyOLCpQsH4k=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/a439e4f37ce71de48d4a979594276e3be0e6278f.patch";
hash = "sha256-X2VPGigK582J+a/u2Kg74w25/+CTpVWU9D3Eqgnb2PU=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/7002dc04863d38c57cfd2e6fc60a1cf5a613fd8e.patch";
relative = "amd/comgr";
})
#[Comgr] Extend ISA compatibility for CCOB
# [Comgr] Extend ISA compatibility for CCOB
(fetchpatch {
sha256 = "sha256-PCi0QHLiEQCTIYRtSSbhOjXANJ3zC3VLdMED1BEfQeg=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/fa80abb77d5ae6f8d89ab956e7ebda9c802a804f.patch";
hash = "sha256-/50I+PqxL3oaQMqg5vR7+ibUcXO1SvfXBdw/sybRt1o=";
url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/2c1e44fc3eacadcafdd4ada3e3184a092b6f26c5.patch";
relative = "amd/comgr";
})
];
postPatch = ''
substituteInPlace cmake/opencl_pch.cmake \
substituteInPlace cmake/opencl_header.cmake \
--replace-fail "\''${CLANG_CMAKE_DIR}/../../../" "${llvm.clang-unwrapped.lib}"
'';
@@ -17,13 +17,13 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rocm-core";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocm-core";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-3WQPR7v7LRoo4ioCXzkUEqnJdXA3uxSgoZZQvEG5Gqg=";
hash = "sha256-AqD6MByFAtY6IzAFXdCrCANPXqROaGHybdKixv3NbXE=";
};
patches = [
@@ -41,6 +41,7 @@ stdenv.mkDerivation (finalAttrs: {
"-DROCM_VERSION=${finalAttrs.version}"
"-DBUILD_ID=${finalAttrs.env.BUILD_ID}"
"-DCMAKE_INSTALL_LIBDIR=lib"
"-DCMAKE_INSTALL_LIBEXECDIR=libexec"
"-DCMAKE_INSTALL_INCLUDEDIR=include"
"-DCMAKE_INSTALL_BINDIR=bin"
];
@@ -37,12 +37,6 @@ stdenv.mkDerivation (finalAttrs: {
patches = [
./cmake.patch
(fetchpatch {
name = "cmake-4-compat-dont-set-cmp0053.patch";
url = "https://github.com/ROCm/llvm-project/commit/a18cc4c7cb51f94182b6018c7c73acde1b8ebddb.patch";
hash = "sha256-kp/Ld0IhjWgRbRR9R/CKdkI9ELvPkQSAMqPsAPFxzhM=";
relative = "amd/device-libs";
})
];
nativeBuildInputs = [
@@ -16,13 +16,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocm-runtime";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "ROCR-Runtime";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-gqe1trGc/Cu1XFA4aYjOzFurUgebLbdTHEJi4iw2+kk=";
hash = "sha256-6xELKQ/uqAoorsCR/H7d8iNK7LsVNsW2DRRZo5cU7UM=";
};
cmakeBuildType = "RelWithDebInfo";
@@ -58,8 +58,8 @@ stdenv.mkDerivation (finalAttrs: {
./queue-failure.patch
(fetchpatch {
# [PATCH] rocr: Extend HIP ISA compatibility check
sha256 = "sha256-8r2Lb5lBfFaZC3knCxfXGcnkzNv6JxOKyJn2rD5gus4=";
url = "https://github.com/GZGavinZhao/rocm-systems/commit/c13cd118fcc8e0bc9ae8de62897542dca7352b71.patch";
hash = "sha256-8r2Lb5lBfFaZC3knCxfXGcnkzNv6JxOKyJn2rD5gus4=";
url = "https://github.com/GZGavinZhao/rocm-systems/commit/dcef23cf896f4dcbc7ed81abeaa4ec2208dcdd8c.patch";
relative = "projects/rocr-runtime";
})
(fetchpatch {
@@ -67,17 +67,6 @@ stdenv.mkDerivation (finalAttrs: {
url = "https://github.com/ROCm/ROCR-Runtime/commit/41bfc66aef437a5b349f71105fa4b907cc7e17d5.patch";
hash = "sha256-A7VhPR3eSsmjq2cTBSjBIz9i//WiNjoXm0EsRKtF+ns=";
})
(fetchpatch {
# [PATCH] hsakmt: bump vgpr count for gfx1151 (#1807) (#1986)
# We apply only the change that adds a GFX_VERSION_GFX1151 define but
# *not* the default vgpr change which causes crashes on old kernels
# This was later partially reverted in # [PATCH] Revert "hsakmt: bump vgpr count for gfx1151 (#1807) (#1986)"
name = "rocr-runtime-gfx1151-vgpr.patch";
url = "https://github.com/ROCm/rocm-systems/commit/09ba45b3f43ec333a84a0ca178fcd1e3ea9400a9.patch";
relative = "projects/rocr-runtime";
includes = [ "libhsakmt/src/libhsakmt.h" ];
hash = "sha256-/V5i+sr88n7fK4yNjR/FpY0ZpiEG5xAD6Oq+9ZOikd4=";
})
(fetchpatch {
# [PATCH] hsakmt: Expose and use CWSR and Control stack sizes (#2200)
# Fixes potential mismatches between stack sizes configured in amdgpu kernel module
@@ -11,13 +11,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocm-smi";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocm_smi_lib";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-/3to1+45n+D8snwrFVhApSMWp/rIXmTFK+BiA0mcKZo=";
hash = "sha256-IeOutA/Mpt/75VpdcfPnZrWvx6GG/JLaQNfBrNKb+Mw=";
};
propagatedBuildInputs = [
@@ -12,14 +12,14 @@
}:
stdenv.mkDerivation (finalAttrs: {
version = "7.1.1";
version = "7.2.0";
pname = "rocminfo";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocminfo";
rev = "rocm-${finalAttrs.version}";
sha256 = "sha256-SuW34m2ep69+dj1rb0vqfQcK83FBtjlMLqSMCUZltU4=";
hash = "sha256-Md91iH2YWU7ziNE//xxKuk2RrXdujO9mY6SZY5R9TqA=";
};
strictDeps = true;
@@ -42,7 +42,7 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rocmlir${suffix}";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -55,7 +55,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocMLIR";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-A9vUvsEZrZlNEW4cscF66L48rJQ1zJYmIzwXQ2QzJ3s=";
hash = "sha256-gg37odxZETpH7HqY+44KQta6f0vxfcxhzMAqxU6vicY=";
};
nativeBuildInputs = [
@@ -15,7 +15,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocprim";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocPRIM";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-/O6aEF23IaUSVzEDRqIlpYpxNCSTa/CwVDgXOhvTrQc=";
hash = "sha256-/6x4S8W4MbWY+M2b46Tmebj2echBIs31OnKq+50M21M=";
};
nativeBuildInputs = [
@@ -18,13 +18,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocprofiler-register";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocprofiler-register";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-XlBz5rDMvnPxyCiqpEpIUWzbHJ1BbMzeCJTrP5kKf+w=";
hash = "sha256-6jr6dfmCjeP5PCNfLDMmh8IO9Hxz6RvLlacNwIWQduQ=";
fetchSubmodules = true;
};
@@ -44,13 +44,13 @@ let
in
stdenv.mkDerivation (finalAttrs: {
pname = "rocprofiler";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocprofiler";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-V+E+6D/mBW/NVdHSBVUVMG5Z/QOF73vfZ5rryW2VBtI=";
hash = "sha256-W32VAdtZdFmkHeUwyHIIHk8yRnABq6gUaHLZ9Bj8kYI=";
fetchSubmodules = true;
};
@@ -12,13 +12,13 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocr-debug-agent";
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rocr_debug_agent";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-zW92HyK4nGcDaJxmBaHeURAivs6dU059WxaywL6dAk4=";
hash = "sha256-w2Zg4kpuKy68DkVGsBTUsjRZoV/Y/Z3Q8s0oSIR3Ask=";
};
nativeBuildInputs = [
@@ -15,7 +15,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocrand${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocRAND";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-QhHxpjr3SqeSSMS4QmnDq8zIxEeUtrgwb9JG84jfJgY=";
hash = "sha256-BX1A74ZBXE0fqYqA7WijezCezUOpEciUm6k2vhFQlA4=";
};
nativeBuildInputs = [
@@ -37,7 +37,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocsolver${clr.gpuArchSuffix}";
version = "7.1.0";
version = "7.2.0";
outputs = [
"out"
@@ -53,7 +53,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocSOLVER";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-U5RrV90UZiIUgLYleLJ6rFHWRFMQNyh3zezRusj7T0M=";
hash = "sha256-aCmNTF8P0Y7piPkRnu/+DWFhTd7X8mDcaQpotfw/4hM=";
};
nativeBuildInputs = [
@@ -7,6 +7,7 @@
cmake,
rocm-cmake,
rocprim,
rocblas,
clr,
gfortran,
gtest,
@@ -19,7 +20,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocsparse${clr.gpuArchSuffix}";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -35,7 +36,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocSPARSE";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-ZNeFslIQZlFCNqWHmrpcxDuOvVu1cjiEd8b6ZNXcZzo=";
hash = "sha256-jmrX1a09DlQTkuXwkrInt9pKX5pk3oCO88S0pKXXhO0=";
};
nativeBuildInputs = [
@@ -48,6 +49,7 @@ stdenv.mkDerivation (finalAttrs: {
buildInputs = [
rocprim
rocblas
]
++ lib.optionals (buildTests || buildBenchmarks) [
gtest
+24 -24
View File
@@ -6,7 +6,7 @@
{
matrix-01 = fetchzip {
sha256 = "sha256-AHur5ZIDZTFRrO2GV0ieXrffq4KUiGWiZ59pv0fUtEQ=";
hash = "sha256-AHur5ZIDZTFRrO2GV0ieXrffq4KUiGWiZ59pv0fUtEQ=";
urls = [
"${mirror1}/SNAP/amazon0312.tar.gz"
@@ -15,7 +15,7 @@
};
matrix-02 = fetchzip {
sha256 = "sha256-0rSxaN4lQcdaCLsvlgicG70FXUxXeERPiEmQ4MzbRdE=";
hash = "sha256-0rSxaN4lQcdaCLsvlgicG70FXUxXeERPiEmQ4MzbRdE=";
urls = [
"${mirror1}/Muite/Chebyshev4.tar.gz"
@@ -24,7 +24,7 @@
};
matrix-03 = fetchzip {
sha256 = "sha256-hDzDWDUnHEyFedX/tMNq83ZH8uWyM4xtZYUUAD3rizo=";
hash = "sha256-hDzDWDUnHEyFedX/tMNq83ZH8uWyM4xtZYUUAD3rizo=";
urls = [
"${mirror1}/FEMLAB/sme3Dc.tar.gz"
@@ -33,7 +33,7 @@
};
matrix-04 = fetchzip {
sha256 = "sha256-GmN2yOt/MoX01rKe05aTyB3ypUP4YbQGOITZ0BqPmC0=";
hash = "sha256-GmN2yOt/MoX01rKe05aTyB3ypUP4YbQGOITZ0BqPmC0=";
urls = [
"${mirror1}/Williams/webbase-1M.tar.gz"
@@ -42,7 +42,7 @@
};
matrix-05 = fetchzip {
sha256 = "sha256-gQNjfVyWzNM9RwImJGhkhahRmZz74LzDs1oijL7mI7k=";
hash = "sha256-gQNjfVyWzNM9RwImJGhkhahRmZz74LzDs1oijL7mI7k=";
urls = [
"${mirror1}/Williams/mac_econ_fwd500.tar.gz"
@@ -51,7 +51,7 @@
};
matrix-06 = fetchzip {
sha256 = "sha256-87cdZjntNcTuz5BtO59irhcuRbPllWSbhCEX3Td02qc=";
hash = "sha256-87cdZjntNcTuz5BtO59irhcuRbPllWSbhCEX3Td02qc=";
urls = [
"${mirror1}/Williams/mc2depi.tar.gz"
@@ -60,7 +60,7 @@
};
matrix-07 = fetchzip {
sha256 = "sha256-WRamuJX3D8Tm+k0q67RjUDG3DeNAxhKiaPkk5afY5eU=";
hash = "sha256-WRamuJX3D8Tm+k0q67RjUDG3DeNAxhKiaPkk5afY5eU=";
urls = [
"${mirror1}/Bova/rma10.tar.gz"
@@ -69,7 +69,7 @@
};
matrix-08 = fetchzip {
sha256 = "sha256-5dhkm293Mc3lzakKxHy5W5XIn4Rw+gihVh7gyrjEHXo=";
hash = "sha256-5dhkm293Mc3lzakKxHy5W5XIn4Rw+gihVh7gyrjEHXo=";
urls = [
"${mirror1}/JGD_BIBD/bibd_22_8.tar.gz"
@@ -78,7 +78,7 @@
};
matrix-09 = fetchzip {
sha256 = "sha256-czjLWCjXAjZCk5TGYHaEkwSAzQu3TQ3QyB6eNKR4G88=";
hash = "sha256-czjLWCjXAjZCk5TGYHaEkwSAzQu3TQ3QyB6eNKR4G88=";
urls = [
"${mirror1}/Hamm/scircuit.tar.gz"
@@ -87,7 +87,7 @@
};
matrix-10 = fetchzip {
sha256 = "sha256-bYuLnJViAIcIejAkh69/bsNAVIDU4wfTLtD+nmHd6FM=";
hash = "sha256-bYuLnJViAIcIejAkh69/bsNAVIDU4wfTLtD+nmHd6FM=";
urls = [
"${mirror1}/Sandia/ASIC_320k.tar.gz"
@@ -96,7 +96,7 @@
};
matrix-11 = fetchzip {
sha256 = "sha256-aDwn8P1khYjo2Agbq5m9ZBInJUxf/knJNvyptt0fak0=";
hash = "sha256-aDwn8P1khYjo2Agbq5m9ZBInJUxf/knJNvyptt0fak0=";
urls = [
"${mirror1}/GHS_psdef/bmwcra_1.tar.gz"
@@ -105,7 +105,7 @@
};
matrix-12 = fetchzip {
sha256 = "sha256-8OJqA/byhlAZd869TPUzZFdsOiwOoRGfKyhM+RMjXoY=";
hash = "sha256-8OJqA/byhlAZd869TPUzZFdsOiwOoRGfKyhM+RMjXoY=";
urls = [
"${mirror1}/HB/nos1.tar.gz"
@@ -114,7 +114,7 @@
};
matrix-13 = fetchzip {
sha256 = "sha256-FS0rKqmg+uHwsM/yGfQLBdd7LH/rUrdutkNGBD/Mh1I=";
hash = "sha256-FS0rKqmg+uHwsM/yGfQLBdd7LH/rUrdutkNGBD/Mh1I=";
urls = [
"${mirror1}/HB/nos2.tar.gz"
@@ -123,7 +123,7 @@
};
matrix-14 = fetchzip {
sha256 = "sha256-DANnlrNJikrI7Pst9vRedtbuxepyHmCIu2yhltc4Qcs=";
hash = "sha256-DANnlrNJikrI7Pst9vRedtbuxepyHmCIu2yhltc4Qcs=";
urls = [
"${mirror1}/HB/nos3.tar.gz"
@@ -132,7 +132,7 @@
};
matrix-15 = fetchzip {
sha256 = "sha256-21mUgqjWGUfYgiWwSrKh9vH8Vdt3xzcefmqYNYRpxiY=";
hash = "sha256-21mUgqjWGUfYgiWwSrKh9vH8Vdt3xzcefmqYNYRpxiY=";
urls = [
"${mirror1}/HB/nos4.tar.gz"
@@ -141,7 +141,7 @@
};
matrix-16 = fetchzip {
sha256 = "sha256-FOuXvGqBBFNkVS6cexmkluret54hCfCOdK+DOZllE4c=";
hash = "sha256-FOuXvGqBBFNkVS6cexmkluret54hCfCOdK+DOZllE4c=";
urls = [
"${mirror1}/HB/nos5.tar.gz"
@@ -150,7 +150,7 @@
};
matrix-17 = fetchzip {
sha256 = "sha256-+7NI1rA/qQxYPpjXKHvAaCZ+LSaAJ4xuJvMRMBEUYxg=";
hash = "sha256-+7NI1rA/qQxYPpjXKHvAaCZ+LSaAJ4xuJvMRMBEUYxg=";
urls = [
"${mirror1}/HB/nos6.tar.gz"
@@ -159,7 +159,7 @@
};
matrix-18 = fetchzip {
sha256 = "sha256-q3NxJjbwGGcFiQ9nhWfUKgZmdVwCfPmgQoqy0AqOsNc=";
hash = "sha256-q3NxJjbwGGcFiQ9nhWfUKgZmdVwCfPmgQoqy0AqOsNc=";
urls = [
"${mirror1}/HB/nos7.tar.gz"
@@ -168,7 +168,7 @@
};
matrix-19 = fetchzip {
sha256 = "sha256-0GAN6qmVfD+tprIigzuUUUwm5KVhkN9X65wMEvFltDY=";
hash = "sha256-0GAN6qmVfD+tprIigzuUUUwm5KVhkN9X65wMEvFltDY=";
urls = [
"${mirror1}/DNVS/shipsec1.tar.gz"
@@ -177,7 +177,7 @@
};
matrix-20 = fetchzip {
sha256 = "sha256-f28Du/Urxsiq5NkRmRO10Zz9vvGRjEchquzHzbZpZ7U=";
hash = "sha256-f28Du/Urxsiq5NkRmRO10Zz9vvGRjEchquzHzbZpZ7U=";
urls = [
"${mirror1}/Cote/mplate.tar.gz"
@@ -186,7 +186,7 @@
};
matrix-21 = fetchzip {
sha256 = "sha256-O+Wy0NfCU1hVUOfNR1dJpvDHLBwwa301IRJDrQJnhak=";
hash = "sha256-O+Wy0NfCU1hVUOfNR1dJpvDHLBwwa301IRJDrQJnhak=";
urls = [
"${mirror1}/Bai/qc2534.tar.gz"
@@ -195,7 +195,7 @@
};
matrix-22 = fetchzip {
sha256 = "sha256-oxMnt8U5Cf1ILWcBdU6W9jdSMMm+U6bIVl8nm3n3+OA=";
hash = "sha256-oxMnt8U5Cf1ILWcBdU6W9jdSMMm+U6bIVl8nm3n3+OA=";
urls = [
"${mirror1}/Chevron/Chevron2.tar.gz"
@@ -204,7 +204,7 @@
};
matrix-23 = fetchzip {
sha256 = "sha256-MFD9BxFI/3IS7yatW121BAI04fbqrXpgYDT5UKjeKcU=";
hash = "sha256-MFD9BxFI/3IS7yatW121BAI04fbqrXpgYDT5UKjeKcU=";
urls = [
"${mirror1}/Chevron/Chevron3.tar.gz"
@@ -213,7 +213,7 @@
};
matrix-24 = fetchzip {
sha256 = "sha256-ikS8O51pe1nt3BNyhvfvqCbVL0+bg/da9bqGqeBDkTg=";
hash = "sha256-ikS8O51pe1nt3BNyhvfvqCbVL0+bg/da9bqGqeBDkTg=";
urls = [
"${mirror1}/Chevron/Chevron4.tar.gz"
@@ -15,7 +15,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocthrust";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocThrust";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-Gncy0wvN8M4JSmnjbxaED+M9rLo4A/7LCCq/6temiZU=";
hash = "sha256-BDxFeCR2FfZI+TGMi2jBDHORCoAnFwk8WGGKMckCkmU=";
};
nativeBuildInputs = [
@@ -19,7 +19,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "roctracer";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -19,7 +19,7 @@
stdenv.mkDerivation (finalAttrs: {
pname = "rocwmma";
version = "7.1.1";
version = "7.2.0";
outputs = [
"out"
@@ -38,7 +38,7 @@ stdenv.mkDerivation (finalAttrs: {
owner = "ROCm";
repo = "rocWMMA";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-6hvdq3J8UhldcyNqXxi9UEF+P0RgG8o5YdIZLzvpgXw=";
hash = "sha256-O+P+pobPnlKnmJQylFp/qUQLE1kJKbxjMWyGS6lc3Mo=";
};
patches = lib.optionals (buildTests || buildBenchmarks) [
@@ -29,13 +29,13 @@ stdenv.mkDerivation (finalAttrs: {
"cpu"
);
version = "7.1.1";
version = "7.2.0";
src = fetchFromGitHub {
owner = "ROCm";
repo = "rpp";
rev = "rocm-${finalAttrs.version}";
hash = "sha256-ZlJRV57ybPsczvoWwd5vr2n7EKje0vj2GJYsvVY+qas=";
hash = "sha256-mET9nZReO3BrNZkcNeB7hOwIQaVZ8OD7lPD3b1Jot2k=";
};
nativeBuildInputs = [
@@ -18,14 +18,14 @@
buildPythonPackage rec {
pname = "tensile";
version = "7.1.1";
version = "7.2.0";
pyproject = true;
src = fetchFromGitHub {
owner = "ROCm";
repo = "Tensile";
rev = "rocm-${version}";
hash = "sha256-uYLlbtXrc8m8/RpMbP7WaS7VYGIUm2LHYzv00DZAzvE=";
hash = "sha256-pTuEeu+0E+RlpdeFho0mJmx/V/tou/ZRLYX9KoWEPpA=";
};
# TODO: It should be possible to run asm caps test ONCE for all supported arches
@@ -58,8 +58,9 @@ buildPythonPackage rec {
./tensile-create-library-dont-copy-twice.diff
(fetchpatch {
# [PATCH] Extend Tensile HIP ISA compatibility
sha256 = "sha256-d+fVf/vz+sxGqJ96vuxe0jRMgbC5K6j5FQ5SJ1e3Sl8=";
url = "https://github.com/GZGavinZhao/Tensile/commit/855cb15839849addb0816a6dde45772034a3e41f.patch";
hash = "sha256-ZHXNYSeLkhhNDaIfdqQm68Pxmh1shUL7mAVmh8/I6Xk=";
url = "https://github.com/GZGavinZhao/rocm-libraries/commit/1f7135dfc0cdb175c8f0e5eb71b2d24699942873.patch";
relative = "shared/tensile";
})
];