Skip to content

Conversation

@YonahGoldberg
Copy link
Contributor

No description provided.

@llvmbot
Copy link
Member

llvmbot commented Dec 19, 2025

@llvm/pr-subscribers-llvm-transforms

Author: Yonah Goldberg (YonahGoldberg)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/173117.diff

1 Files Affected:

  • (modified) llvm/lib/Transforms/Scalar/SROA.cpp (+16-3)
diff --git a/llvm/lib/Transforms/Scalar/SROA.cpp b/llvm/lib/Transforms/Scalar/SROA.cpp
index f2a85cc7af441..0673360abaef3 100644
--- a/llvm/lib/Transforms/Scalar/SROA.cpp
+++ b/llvm/lib/Transforms/Scalar/SROA.cpp
@@ -1523,12 +1523,14 @@ LLVM_DUMP_METHOD void AllocaSlices::dump() const { print(dbgs()); }
 
 /// Walk the range of a partitioning looking for a common type to cover this
 /// sequence of slices.
-static std::pair<Type *, IntegerType *>
+/// Returns: {CommonType, LargestIntegerType, OnlyIntrinsicUsers}
+static std::tuple<Type *, IntegerType *, bool>
 findCommonType(AllocaSlices::const_iterator B, AllocaSlices::const_iterator E,
                uint64_t EndOffset) {
   Type *Ty = nullptr;
   bool TyIsCommon = true;
   IntegerType *ITy = nullptr;
+  bool OnlyIntrinsicUsers = true;
 
   // Note that we need to look at *every* alloca slice's Use to ensure we
   // always get consistent results regardless of the order of slices.
@@ -1536,6 +1538,8 @@ findCommonType(AllocaSlices::const_iterator B, AllocaSlices::const_iterator E,
     Use *U = I->getUse();
     if (isa<IntrinsicInst>(*U->getUser()))
       continue;
+    // We found a non-intrinsic user
+    OnlyIntrinsicUsers = false;
     if (I->beginOffset() != B->beginOffset() || I->endOffset() != EndOffset)
       continue;
 
@@ -1569,7 +1573,7 @@ findCommonType(AllocaSlices::const_iterator B, AllocaSlices::const_iterator E,
       Ty = UserTy;
   }
 
-  return {TyIsCommon ? Ty : nullptr, ITy};
+  return {TyIsCommon ? Ty : nullptr, ITy, OnlyIntrinsicUsers};
 }
 
 /// PHI instructions that use an alloca and are subsequently loaded can be
@@ -5257,7 +5261,7 @@ selectPartitionType(Partition &P, const DataLayout &DL, AllocaInst &AI,
 
   // Check if there is a common type that all slices of the partition use that
   // spans the partition.
-  auto [CommonUseTy, LargestIntTy] =
+  auto [CommonUseTy, LargestIntTy, OnlyIntrinsicUsers] =
       findCommonType(P.begin(), P.end(), P.endOffset());
   if (CommonUseTy) {
     TypeSize CommonUseSize = DL.getTypeAllocSize(CommonUseTy);
@@ -5295,6 +5299,15 @@ selectPartitionType(Partition &P, const DataLayout &DL, AllocaInst &AI,
         isIntegerWideningViable(P, LargestIntTy, DL))
       return {LargestIntTy, true, nullptr};
 
+    // If there are only intrinsic users of an aggregate type, try to
+    // represent as a legal integer type because we are probably just copying
+    // data around and the integer can be promoted.
+    if (OnlyIntrinsicUsers && DL.isLegalInteger(P.size() * 8) &&
+        TypePartitionTy->isAggregateType())
+      auto *IntNTy = Type::getIntNTy(*C, P.size() * 8);
+      return {IntNTy, isIntegerWideningViable(P, IntNTy, DL), nullptr};
+    }
+
     // Fallback to TypePartitionTy and we probably won't promote.
     return {TypePartitionTy, false, nullptr};
   }

@github-actions
Copy link

github-actions bot commented Dec 19, 2025

✅ With the latest revision this PR passed the C/C++ code formatter.

@github-actions
Copy link

github-actions bot commented Dec 19, 2025

🪟 Windows x64 Test Results

  • 128820 tests passed
  • 2836 tests skipped
  • 1 test failed

Failed Tests

(click on a test name to see its output)

LLVM

LLVM.CodeGen/NVPTX/lower-byval-args.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
c:\_work\llvm-project\llvm-project\build\bin\opt.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\opt.exe' -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 3
c:\_work\llvm-project\llvm-project\build\bin\opt.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\opt.exe' -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 4
c:\_work\llvm-project\llvm-project\build\bin\opt.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\opt.exe' -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 5
c:\_work\llvm-project\llvm-project\build\bin\opt.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\opt.exe' -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 6
c:\_work\llvm-project\llvm-project\build\bin\opt.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=COMMON,COPY
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\opt.exe' -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=COMMON,COPY
# note: command had no output on stdout or stderr
# RUN: at line 7
c:\_work\llvm-project\llvm-project\build\bin\llc.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll -mcpu=sm_60 -mattr=ptx77 | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll --check-prefixes=PTX,PTX_60
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -mcpu=sm_60 -mattr=ptx77
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll' --check-prefixes=PTX,PTX_60
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll:458:13: error: PTX-NEXT: expected string not found in input
# | ; PTX-NEXT: .reg .b32 %r<3>;
# |             ^
# | <stdin>:295:17: note: scanning from here
# |  .reg .b64 %SPL;
# |                 ^
# | <stdin>:296:2: note: possible intended match here
# |  .reg .b64 %rd<30>;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\NVPTX\lower-byval-args.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |           290:  .param .align 4 .b8 memcpy_to_param_param_1[8] 
# |           291: ) // @memcpy_to_param 
# |           292: { 
# |           293:  .local .align 8 .b8 __local_depot9[8]; 
# |           294:  .reg .b64 %SP; 
# |           295:  .reg .b64 %SPL; 
# | next:458'0                     X error: no match found
# |           296:  .reg .b64 %rd<30>; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~
# | next:458'1      ?                   possible intended match
# |           297:  
# | next:458'0     ~
# |           298: // %bb.0: // %entry 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~
# |           299:  mov.b64 %SPL, __local_depot9; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           300:  cvta.local.u64 %SP, %SPL; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           301:  ld.param.b64 %rd1, [memcpy_to_param_param_0]; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

If these failures are unrelated to your changes (for example tests are broken or flaky at HEAD), please open an issue at https://github.com/llvm/llvm-project/issues and add the infrastructure label.

@github-actions
Copy link

github-actions bot commented Dec 19, 2025

🐧 Linux x64 Test Results

  • 167340 tests passed
  • 2968 tests skipped
  • 1 test failed

Failed Tests

(click on a test name to see its output)

LLVM

LLVM.CodeGen/NVPTX/lower-byval-args.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 3
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 4
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 5
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 6
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,COPY
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,COPY
# note: command had no output on stdout or stderr
# RUN: at line 7
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mcpu=sm_60 -mattr=ptx77 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=PTX,PTX_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mcpu=sm_60 -mattr=ptx77
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=PTX,PTX_60
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll:458:13: error: PTX-NEXT: expected string not found in input
# | ; PTX-NEXT: .reg .b32 %r<3>;
# |             ^
# | <stdin>:295:17: note: scanning from here
# |  .reg .b64 %SPL;
# |                 ^
# | <stdin>:296:2: note: possible intended match here
# |  .reg .b64 %rd<30>;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |           290:  .param .align 4 .b8 memcpy_to_param_param_1[8] 
# |           291: ) // @memcpy_to_param 
# |           292: { 
# |           293:  .local .align 8 .b8 __local_depot9[8]; 
# |           294:  .reg .b64 %SP; 
# |           295:  .reg .b64 %SPL; 
# | next:458'0                     X error: no match found
# |           296:  .reg .b64 %rd<30>; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~
# | next:458'1      ?                   possible intended match
# |           297:  
# | next:458'0     ~
# |           298: // %bb.0: // %entry 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~
# |           299:  mov.b64 %SPL, __local_depot9; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           300:  cvta.local.u64 %SP, %SPL; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           301:  ld.param.b64 %rd1, [memcpy_to_param_param_0]; 
# | next:458'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

If these failures are unrelated to your changes (for example tests are broken or flaky at HEAD), please open an issue at https://github.com/llvm/llvm-project/issues and add the infrastructure label.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants