Port publication and performance utility fixes

README.md

@@ -107,9 +107,7 @@ See [SECURITY.md](SECURITY.md) for our responsible disclosure policy.
 
 ## Contact
 
-For questions, issues, or contributions, please reach out to the maintainers:
-
-- Chris Austen — [@causten](https://github.com/causten) · chausten@amd.com
+For questions, [open a GitHub issue](../../issues/new).
 
 See [CODEOWNERS](.github/CODEOWNERS) for the full ownership list.
 

docs/PR_REVIEW_CHECKLIST.md

@@ -32,6 +32,11 @@ the rationale.
 
 - Unreleased hardware codenames, unannounced chip IDs, or NDA features in
   code, comments, commits, or docs.
+- Internal-only hyperlinks or URLs in code, comments, commits, docs, test
+  data, or metadata, including internal Jira/OnTrack, Confluence,
+  source-control repositories, and internal network hosts. Keep non-sensitive
+  ticket identifiers as plain text when provenance is useful, but never
+  include an internal URL.
 - C++ exceptions (`throw`, `try`/`catch`); use `LogicalResult` /
   `emitOpError` / `signalPassFailure` instead.
 - RTTI (`dynamic_cast`, `typeid`); use LLVM's `isa`/`cast`/`dyn_cast`.

mlir/include/mlir/Dialect/Rock/IR/RockAttrDefs.td

@@ -600,7 +600,7 @@ def Rock_BlockwiseMatrixParamsAttr : Rock_Attr<"BlockwiseMatrixParams", []> {
     - swapThreadIterSubDims: Trick to reduce LDS bank conflicts (see more info here: https://github.com/ROCm/rocMLIR/pull/1209)
     - LDSLayoutDxK: Wheter the layout in LDS is DxK
     - directToLDS: Wheter direct to LDS is enabled
-    - splitKAcrossThreadsFirst: Used for attention, when bypassing LDS for the result of the first GEMM, explanation here: https://github.com/ROCm/rocMLIR-internal/issues/1201#issuecomment-1898925539
+    - splitKAcrossThreadsFirst: Used for attention when bypassing LDS for the result of the first GEMM (ROCm/rocMLIR-internal#1201).
     - g: gemm parameter G
     - d: gemm parameter D (could be M or N)
     - inDPerThread: How many elements of D (M or N) each thread is going to load from memory.

mlir/lib/Dialect/Rock/Transforms/GridwiseGemmToBlockwise.cpp

@@ -2329,11 +2329,10 @@ struct GridwiseAttentionAccelRewritePattern
       ldsLayoutCfgNG0.doSwapThreadIterSubDims = false;
     }
     if (op.getEnableSoftmax()) {
-      // TODO: Workaround for issue
-      // https://github.com/ROCm/rocMLIR-internal/issues/1802 If sumRowBuffer
-      // and expMaxDiffRowBuffer are filled with doSwapThreadIterSubDims=true,
-      // it does not match with the second GEMM N dimension. Find a good
-      // solution to this.
+      // TODO(ROCm/rocMLIR-internal#1802): If sumRowBuffer and
+      // expMaxDiffRowBuffer are filled with doSwapThreadIterSubDims=true, it
+      // does not match with the second GEMM N dimension. Find a good solution
+      // to this.
       ldsLayoutCfgNG0.doSwapThreadIterSubDims = false;
     }
     FailureOr<VectorDimInfo> maybeVectorDimInfoK =

mlir/test/e2e/conv_regression_bwd.toml

@@ -52,11 +52,9 @@ config = "-groupsize=1 -batchsize=64 -in_channels=256 -out_channels=256 -in_h=14
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=64 -out_channels=64 -in_h=4 -in_w=4 -fil_h=2 -fil_w=2 -dilation_h=1 -dilation_w=1 -conv_stride_h=2 -conv_stride_w=2 -padding_h_l=2 -padding_h_r=1 -padding_w_l=2 -padding_w_r=0"
 
-## The following configs are reported from various tickets
+## Regression configurations
 
-############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/70 #
-############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#70
 [[suite.test]]
 config = "-groupsize=1 -batchsize=16 -in_channels=32 -out_channels=32 -in_h=14 -in_w=14 -fil_h=1 -fil_w=1 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
@@ -78,9 +76,7 @@ config = "-groupsize=1 -batchsize=64 -in_channels=32 -out_channels=32 -in_h=14 -
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=64 -out_channels=64 -in_h=14 -in_w=14 -fil_h=1 -fil_w=1 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/127 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#127
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=256 -out_channels=128 -in_h=28 -in_w=28 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
@@ -90,18 +86,14 @@ config = "-groupsize=1 -batchsize=512 -in_channels=256 -out_channels=512 -in_h=7
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=256 -out_channels=64 -in_h=56 -in_w=56 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
-############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/71 #
-############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#71
 [[suite.test]]
 config = "-groupsize=1 -batchsize=32 -in_channels=32 -out_channels=32 -in_h=7 -in_w=7 -fil_h=1 -fil_w=1 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=32 -out_channels=32 -in_h=7 -in_w=7 -fil_h=1 -fil_w=1 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/136 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#136
 [[suite.test]]
 config = "-groupsize=1 -batchsize=256 -in_channels=32 -out_channels=32 -in_h=28 -in_w=28 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=2 -conv_stride_w=2 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 

mlir/test/e2e/conv_regression_fwd.toml

@@ -135,11 +135,9 @@ config = "-groupsize=1 -batchsize=256 -in_channels=64 -out_channels=64 -in_h=56
 [[suite.test]]
 config = "-groupsize=1 -batchsize=256 -in_channels=64 -out_channels=64 -in_h=56 -in_w=56 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
-## The following configs are reported from various tickets
+## Regression configurations
 
-############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/41 #
-############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#41
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=16 -in_w=16 -fil_h=3 -fil_w=3 -dilation_h=2 -dilation_w=2 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
@@ -176,9 +174,7 @@ config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32 -in_w=32 -fil_h=5 -fil_w=5 -dilation_h=2 -dilation_w=2 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
-############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/40 #
-############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#40
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=64 -in_h=32 -in_w=32 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
@@ -194,9 +190,7 @@ config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=8 -out_channels=128 -in_h=16 -in_w=64 -fil_h=3 -fil_w=5 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/114 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#114
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32 -in_w=32 -fil_h=1 -fil_w=1 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
@@ -209,17 +203,13 @@ config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=8 -out_channels=128 -in_h=32 -in_w=32 -fil_h=1 -fil_w=1 -dilation_h=2 -dilation_w=2 -conv_stride_h=2 -conv_stride_w=2 -padding_h_l=0 -padding_h_r=0 -padding_w_l=0 -padding_w_r=0"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/136 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#136
 [[suite.test]]
 config = "-groupsize=1 -batchsize=32 -in_channels=1 -out_channels=64 -in_h=14 -in_w=14 -fil_h=14 -fil_w=14 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
 [[suite.test]]
 config = "-groupsize=1 -batchsize=32 -in_channels=1 -out_channels=32 -in_h=14 -in_w=14 -fil_h=14 -fil_w=14 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/155 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#155
 [[suite.test]]
 config = "-groupsize=1 -batchsize=64 -in_channels=4 -out_channels=64 -in_h=4 -in_w=4 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"

mlir/test/e2e/conv_regression_fwd_navi3x.toml

@@ -33,9 +33,7 @@ config = "-p -rand_side filter"
 [[suite.test]]
 config = "-p -rand_side input"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/127 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#127
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=256 -out_channels=128 -in_h=28 -in_w=28 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 

mlir/test/e2e/conv_regression_fwd_nonNavi3x.toml

@@ -33,9 +33,7 @@ config = "-p -rand_side filter"
 [[suite.test]]
 config = "-p -rand_side input"
 
-#############################################################################################
-# Cases reported in https://github.com/ROCm/rocMLIR-internal/issues/127 #
-#############################################################################################
+# Cases reported in ROCm/rocMLIR-internal#127
 [[suite.test]]
 config = "-groupsize=1 -batchsize=128 -in_channels=256 -out_channels=128 -in_h=28 -in_w=28 -fil_h=3 -fil_w=3 -dilation_h=1 -dilation_w=1 -conv_stride_h=1 -conv_stride_w=1 -padding_h_l=1 -padding_h_r=1 -padding_w_l=1 -padding_w_r=1"
 

mlir/test/fusion/linalg-generic-const-initializer.mlir

@@ -16,7 +16,7 @@
 #transform_map27 = #rock.transform_map<#map24 by [<AddDim{1} ["exp0"] at [0] -> [] at []>, <AddDim{1} ["exp1"] at [1] -> [] at []>, <AddDim{1} ["exp2"] at [2] -> [] at []>] bounds = [1, 1, 1] -> []>
 #transform_map28 = #rock.transform_map<#map25 by [<Broadcast{1} ["dim0"] at [0] -> ["dim0"] at [0]>, <Broadcast{1} ["dim1"] at [1] -> ["dim1"] at [1]>, <Broadcast{1} ["dim2"] at [2] -> ["dim2"] at [2]>] bounds = [32, 384, 3072] -> [1, 1, 1]>
 // A cut down version of the input from
-// https://github.com/ROCm/rocMLIR-internal/issues/1098
+// ROCm/rocMLIR-internal#1098
 // right before it headed down to linalg.generic. The actuall gemm part has been
 // removed for test simplicity.
 module {

mlir/test/fusion/nightly-misc-e2e/issue-940.mlir

@@ -1,4 +1,4 @@
-// The test case that was used to reproduce https://github.com/ROCm/rocMLIR-internal/issues/940
+// Reproduces ROCm/rocMLIR-internal#940.
 // RUN: rocmlir-gen -fut mlir_dot --arch %arch --clone-harness %s | rocmlir-driver -kernel-pipeline=migraphx,highlevel -host-pipeline=migraphx,highlevel | rocmlir-gen -ph -print-results -rand 1 -rand_type float -fut mlir_dot_wrapper --verifier clone - | rocmlir-driver -host-pipeline mhal,runner -kernel-pipeline full -targets %arch | xmir-runner --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext,%conv_validation_wrapper_library_dir/libconv-validation-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext,%linalg_test_lib_dir/libmlir_float16_utils%shlibext,%linalg_test_lib_dir/libmlir_c_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 

mlir/test/fusion/pr-e2e/attention/mixr-attention-small-decode.mlir

@@ -1,5 +1,5 @@
-// This is a design that was in the MIGraphX CI that was previously failing
-// here: https://ontrack-internal.amd.com/browse/SWDEV-558297
+// Regression test for a small attention decode design from the MIGraphX CI
+// (SWDEV-558297).
 
 // RUN: rocmlir-gen -fut mlir_attention --arch %arch --clone-harness %s | rocmlir-driver -kernel-pipeline=migraphx,highlevel -host-pipeline=migraphx,highlevel | rocmlir-gen -ph -rand 1 -rand_type float -fut mlir_attention_wrapper --verifier clone - | rocmlir-driver -host-pipeline mhal -kernel-pipeline full | xmir-runner --shared-libs=%linalg_test_lib_dir/libmlir_rocm_runtime%shlibext,%conv_validation_wrapper_library_dir/libconv-validation-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext,%linalg_test_lib_dir/libmlir_float16_utils%shlibext,%linalg_test_lib_dir/libmlir_c_runner_utils%shlibext,%linalg_test_lib_dir/libmlir_async_runtime%shlibext --entry-point-result=void | FileCheck %s
 

mlir/test/fusion/pr-e2e/gemm-layouts/other/broadcasted-b-e2e.mlir

@@ -12,7 +12,7 @@
 // VECTORIZATION-NEXT: bVectorLen: 8
 
 // XFAIL: *
-// COM: TODO: Fails due to a bug in FoldBroadcast: https://github.com/ROCm/rocMLIR-internal/issues/1746
+// COM: TODO(ROCm/rocMLIR-internal#1746): FoldBroadcast fails for this case.
 
 module {
   func.func @test(%arg0: !migraphx.shaped<1x320x4096xf16, 1310720x1x320>, %arg1: !migraphx.shaped<1x640x320xf16, 204800x1x640>, %arg2: !migraphx.shaped<2x64x10x64x64xf16, 0x10x1x40960x640>) -> !migraphx.shaped<2x64x10x64x64xf16, 2621440x10x1x40960x640> {

mlir/tools/rocmlir-lib/CMakeLists.txt

@@ -266,7 +266,7 @@ if(BUILD_FAT_LIBROCKCOMPILER)
   rocm_create_package(
     NAME ${CMAKE_PROJECT_NAME}
     DESCRIPTION "MLIR packages"
-    MAINTAINER "rocMLIR Dev Team dl.dl-mlir@amd.com"
+    MAINTAINER "ROCm Dev Support <rocm-dev.support@amd.com>"
     HEADER_ONLY
   )
 endif()

mlir/utils/jenkins/Dockerfile

@@ -1,5 +1,5 @@
 FROM ubuntu:24.04
-MAINTAINER Christopher Austen <chausten@amd.com>
+LABEL maintainer="ROCm Dev Support <rocm-dev.support@amd.com>"
 
 ARG ROCM_VERSION=7.2.4
 ARG ROCM_DEB_REPO=http://repo.radeon.com/rocm/apt/${ROCM_VERSION}

mlir/utils/performance/common/CMakeLists.txt

@@ -2,7 +2,10 @@ find_package(hip PATHS /opt/rocm)
 
 if (hip_FOUND)
   add_library(benchmark-driver-utils EXCLUDE_FROM_ALL benchmarkUtils.cpp)
-  target_link_libraries(benchmark-driver-utils PUBLIC hip::host hip::amdhip64)
+  target_link_libraries(benchmark-driver-utils
+    PUBLIC hip::host hip::amdhip64
+    PRIVATE LLVMSupport
+  )
   set_target_properties(benchmark-driver-utils
     PROPERTIES
     LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib"

mlir/utils/performance/common/benchmarkUtils.cpp

@@ -11,6 +11,8 @@
 #include "benchmarkUtils.h"
 #include "hip_f8_impl.h"
 
+#include "llvm/ADT/APFloat.h"
+
 #include <algorithm>
 #include <cassert>
 #include <cstdint>
@@ -56,33 +58,12 @@ uint16_t float_to_bfloat16(float src_val) {
   return target_val.ushortvec[1];
 }
 
-// F16 conversion (does not support Inf or NaN)
-// Reference-1: https://stackoverflow.com/a/1659563/4066096
-// Reference-2: https://arxiv.org/pdf/2112.08926.pdf (page 28)
-uint16_t float_to_float16(float flt) {
-  union {
-    float f;
-    uint32_t u;
-  } x{flt};
-
-  const uint32_t b = x.u + 0x00001000;          // round-to-nearest-even
-  const uint32_t e = (b & 0x7F800000) >> 23;    // exponent
-  const uint32_t m = b & 0x007FFFFF;            // mantissa
-  const uint32_t sign = (b & 0x80000000) >> 16; // sign
-
-  if (e > 112)
-    // normalized case
-    return sign | (((e - 112) << 10) & 0x7C00) | m >> 13;
-
-  if ((e > 101) && (e < 113))
-    // denormalized case
-    return sign | ((((0x007FF000 + m) >> (125 - e)) + 1) >> 1);
-
-  if (e > 143)
-    // saturate
-    return 0x7FFF;
-
-  return sign;
+uint16_t float_to_float16(float value) {
+  llvm::APFloat converted(value);
+  bool losesInfo;
+  converted.convert(llvm::APFloat::IEEEhalf(),
+                    llvm::APFloat::rmNearestTiesToEven, &losesInfo);
+  return static_cast<uint16_t>(converted.bitcastToAPInt().getZExtValue());
 }
 
 // Check if device uses FNUZ FP8 format
@@ -431,12 +412,12 @@ void *makeHostConstant(float flt, DataType computeDataType) {
   }
   case DataType::F16: {
     uint16_t *ret = reinterpret_cast<uint16_t *>(malloc(2));
-    *ret = float_to_bfloat16(flt);
+    *ret = float_to_float16(flt);
     return ret;
   }
   case DataType::BF16: {
     uint16_t *ret = reinterpret_cast<uint16_t *>(malloc(2));
-    *ret = float_to_float16(flt);
+    *ret = float_to_bfloat16(flt);
     return ret;
   }
   case DataType::I8: {

mlir/utils/performance/reportUtils.py

@@ -2,7 +2,7 @@
 import pandas as pd
 import scipy.stats
 
-from typing import Tuple, List
+from typing import List, Tuple
 
 PERF_REPORT_FILE = {
     'hipBLASLt': 'mlir_vs_hipblaslt_perf.csv',
@@ -115,11 +115,16 @@ def set_common_styles(styler: 'pd.io.formats.style.Styler', speedup_cols: list,
             styler.map(colorizer, subset=[col])
 
 
-# Adapted from
-# https://stackoverflow.com/questions/54405704/check-if-all-values-in-dataframe-column-are-the-same
-def unique_cols(df: pd.DataFrame) -> List[str]:
-    a: np.array = df.to_numpy()
-    return df.columns[(a[0] == a).all(0)]
+def constant_columns(data: pd.DataFrame) -> List[str]:
+    """Return columns containing exactly one distinct value and no missing values.
+
+    An entirely missing column represents absent data rather than a constant
+    value and must remain visible in the report.
+    """
+    with_na = data.nunique(dropna=False)
+    without_na = data.nunique(dropna=True)
+    constant_mask = with_na.eq(1) & without_na.eq(1)
+    return with_na.index[constant_mask].to_list()
 
 
 def clean_data_for_humans(data: pd.DataFrame, title: str)\
@@ -147,7 +152,7 @@ def clean_data_for_humans(data: pd.DataFrame, title: str)\
             data.rename(columns={"InputLayout": "Layout"}, inplace=True)
             index_cols["InputLayout"] = "Layout"
 
-    columns_to_drop = unique_cols(data)
+    columns_to_drop = constant_columns(data)
     # Do not drop unique columns in attention for now
     # to keep it transparent what we are tracking.
     # We can revisit this if it ever becomes an issue.