[NVPTX] Enhance vectorization of ld.param & st.param

[kovdan01]

Summary

Since function parameters and return values are passed via param space, we
can force special alignment for values hold in it which will add vectorization
options. This change may be done if the function has private or internal
linkage. Special alignment is forced during 2 phases.

1. Instruction selection lowering. Here we use special alignment for function prototypes (changing both own return value and parameters alignment), call lowering (changing both callee's return value and parameters alignment).

2. IR pass nvptx-lower-args. Here we change alignment of byval parameters that belong to param space (or are casted to it). We only handle cases when all uses of such parameters are loads from it. For such loads, we can change the alignment according to special type alignment and the load offset. Then, load-store-vectorizer IR pass will perform vectorization where alignment allows it.

Special alignment calculated as maximum from default ABI type alignment and
alignment 16. Alignment 16 is chosen because it's the maximum size of
vectorized ld.param & st.param.

Before specifying such special alignment, we should check if it is a multiple
of the alignment that the type already has. For example, if a value has an
enforced alignment of 64, default ABI alignment of 4 and special alignment
of 16, we should preserve 64.

This patch will be followed by a refactoring patch that removes duplicating
code in handling byval and non-byval arguments.

[tra]

FYI. I've recently proposed to pass the values directly, instead of a byval pointer: https://discourse.llvm.org/t/nvptx-calling-convention-for-aggregate-arguments-passed-by-value/5881
It needs additional changes to get lowering done correctly. I should have them ready in about a week.

This change may be done if the function has private or internal linkage.

I think we should be able to do that to all no-kernel functions if we're compiling without -fgpu-rdc. I think we do reduce visibility of non-kernels in that case, but it would be good to make sure.

If S is a multiple of 4 * A, let special alignment be 4 * A.
Else, if S is a multiple of 2 * A, let special alignment be 2 * A.
Else, let special alignment be A.

I'm not sure if that logic makes sense to me. E.g. if we have [5 x i32], the optimal way to load it would be to align it by 16 and use ld.v4 + ld. With your approach [4 x i32] would be loaded ad ld.v4, [6 x i32] as 3 * ld.v2, but [5 x i32] would use 5 * ld. If we can align 4 and 6 element arrays, I do not see why we would not be allowed to align 5-element array, too -- it's an equivalent of struct { [4 x i32], i32 } as far as in-memory layout is concerned.

I wonder if we can just always set alignment to 16. Byval pointer for NVPTX is a fiction anyways as we always copy the data when we actually lower those arguments.

[kovdan01]

In D120129#3338254, @tra wrote:

FYI. I've recently proposed to pass the values directly, instead of a byval pointer: https://discourse.llvm.org/t/nvptx-calling-convention-for-aggregate-arguments-passed-by-value/5881
It needs additional changes to get lowering done correctly. I should have them ready in about a week.

Okay, please let me know when some results are present. Anyway, IMO, we should be able to handle IR with byval pointers used - at least because a person might potentially use frontend different from clang.

This change may be done if the function has private or internal linkage.

I think we should be able to do that to all no-kernel functions if we're compiling without -fgpu-rdc. I think we do reduce visibility of non-kernels in that case, but it would be good to make sure.

Okay, I'll investigate it.

If S is a multiple of 4 * A, let special alignment be 4 * A.
Else, if S is a multiple of 2 * A, let special alignment be 2 * A.
Else, let special alignment be A.

I'm not sure if that logic makes sense to me. E.g. if we have [5 x i32], the optimal way to load it would be to align it by 16 and use ld.v4 + ld. With your approach [4 x i32] would be loaded ad ld.v4, [6 x i32] as 3 * ld.v2, but [5 x i32] would use 5 * ld. If we can align 4 and 6 element arrays, I do not see why we would not be allowed to align 5-element array, too -- it's an equivalent of struct { [4 x i32], i32 } as far as in-memory layout is concerned.

I wonder if we can just always set alignment to 16. Byval pointer for NVPTX is a fiction anyways as we always copy the data when we actually lower those arguments.

I suppose that you are correct and we can always set alignment to 16. The reason why I implemented such logic is that I try to be as conservative as possible. For example, if we have two values of [5 x i32] aligned as 4, they might be placed together without any gaps. If we align them as 16 - an additional gap of 12 bytes will appear. If keeping such king of "layout" stable is not important - I'll change the logic so alignment is always 16 in param space.

[tra]

In D120129#3350839, @kovdan01 wrote:

I suppose that you are correct and we can always set alignment to 16. The reason why I implemented such logic is that I try to be as conservative as possible.

For example, if we have two values of [5 x i32] aligned as 4, they might be placed together without any gaps.

I do not think PTX gives us any guarantees on how parameters are actually arranged. For all we know they may be packed in registers and alignment is completely irrelevant.
I've poked a bit at the SASS generated for various args and alignments and it appears that first few arguments are actually passed via registers. Up to 12 ints are passed via registers on sm_86, Subsequent arguments are passed via local memory.
Here's an example of alignment effect on the argument layout: https://cuda.godbolt.org/z/87Mj1hz4h

In any case, if the code ends up with local memory loads/stores in the hot path, we can already forget about performance and being able to use vectorized local/stores will not help much.
I think the best we can do here is give ptxas as much flexibility as we can. Setting alignment for return values to 16 should be safe. We rarely return multiple values.

If we align them as 16 - an additional gap of 12 bytes will appear. If keeping such king of "layout" stable is not important - I'll change the logic so alignment is always 16 in param space.

Larger alignment will potentially waste some of the limited param space, so it's a trade-off between guaranteeing efficient loads/stores of parameters vs not being able to pass very large number of arguments. E.g. if someone were to have a function with hundreds of [5 x i32] arguments, we would probably be able to handle more such args if they were aligned by 4.
Efficient loads/stores are arguably much more important as they are very common and affect about everyone.

[kovdan01]

@tra Thanks for your comments! Updated the patch according the discussion about forcing alignment 16.

I think we should be able to do that to all no-kernel functions if we're compiling without -fgpu-rdc. I think we do reduce visibility of non-kernels in that case, but it would be good to make sure.

Checked if we do reduce visibility in such cases, and looks like we do not. The following code:

__device__ int foo(int a, int b, int c) {
  return (a + b) / c;
}

Compiles to the following IR:

; ModuleID = 'device.cu'
source_filename = "device.cu"
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

; Function Attrs: convergent mustprogress noinline nounwind optnone
define dso_local noundef i32 @_Z3fooiii(i32 noundef %a, i32 noundef %b, i32 noundef %c) #0 {
entry:
  %a.addr = alloca i32, align 4
  %b.addr = alloca i32, align 4
  %c.addr = alloca i32, align 4
  store i32 %a, i32* %a.addr, align 4
  store i32 %b, i32* %b.addr, align 4
  store i32 %c, i32* %c.addr, align 4
  %0 = load i32, i32* %a.addr, align 4
  %1 = load i32, i32* %b.addr, align 4
  %add = add nsw i32 %0, %1
  %2 = load i32, i32* %c.addr, align 4
  %div = sdiv i32 %add, %2
  ret i32 %div
}

attributes #0 = { convergent mustprogress noinline nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_35" "target-features"="+ptx75,+sm_35" }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4, !5}

!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 11, i32 5]}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 4, !"nvvm-reflect-ftz", i32 0}
!3 = !{i32 7, !"frame-pointer", i32 2}
!4 = !{!"clang version 15.0.0 (https://github.com/llvm/llvm-project.git 9879c555f21097aee15e73dd25bd89f652dba8ea)"}
!5 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}

The compilation command that I use:

clang++ --cuda-gpu-arch=sm_35 --cuda-path=/opt/cuda/ -S -emit-llvm -fno-gpu-rdc device.cu

Also, searching for GPURelocatableDeviceCode through LLVM codebase does not get results where this value is checked in context of reducing function visibility. I could implement that change, and IMHO that should be a separate patch.

Regarding the current patch – how is https://reviews.llvm.org/D118084 going? Can we merge this patch without waiting for your change about passing byval aggregates directly? As I already mentioned in my previous comment, I suppose the change is useful as far as compiler should work good with any type of IR maybe even generated by non-clang frontend.

[tra]

In D120129#3390733, @kovdan01 wrote:

@tra Thanks for your comments! Updated the patch according the discussion about forcing alignment 16.

I think we should be able to do that to all no-kernel functions if we're compiling without -fgpu-rdc. I think we do reduce visibility of non-kernels in that case, but it would be good to make sure.

Checked if we do reduce visibility in such cases, and looks like we do not.

I was indeed mistaken. AFAICT, we only internalize some device-side variables.
@yaxunl - Sam, do we change visibility for anything else? I know we must keep kernels visible, but the question is whether we ever internalize non-kernel functions and if not, whether we want to. In this case it would allow us to bump argument and return value alignment.

Also, searching for GPURelocatableDeviceCode through LLVM codebase does not get results where this value is checked in context of reducing function visibility. I could implement that change, and IMHO that should be a separate patch.

Yeah, it's a separate issue.

Regarding the current patch – how is https://reviews.llvm.org/D118084 going? Can we merge this patch without waiting for your change about passing byval aggregates directly? As I already mentioned in my previous comment, I suppose the change is useful as far as compiler should work good with any type of IR maybe even generated by non-clang frontend.

There's been little progress on that patch lately. I've got way less time to work on all related changes than I hoped for. This patch does not depend on my changes and will not be blocked. I was mostly pointing at it as a heads up that if/when it lands we would no longer see byval pointers in clang-generated IR and the impact of this patch would be reduced.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

4319–4320

"..may want to increase their alignment in a way that ensures that we can effectively vectorize their loads & stores."

4321

"or has private linkage"

4322

Did you mean "rely on default alignment" instead of "linkage"?

"To allow using 128-bit vectorized loads/stores. his function ensures that alignment is 16 or greater."

llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp

235

"tryIncreaseByValAlign" does not sound quite right.
"adjustByValArgAlignment()" ?

I'd also mention that we traverse all loads from byval pointer and adjust their alignment, if those were using known offset.

llvm/test/CodeGen/NVPTX/param-vectorize.ll

39

We should also check that alignment of kernel parameters is not affected.

[yaxunl]

In D120129#3390959, @tra wrote:

In D120129#3390733, @kovdan01 wrote:

@tra Thanks for your comments! Updated the patch according the discussion about forcing alignment 16.

I think we should be able to do that to all no-kernel functions if we're compiling without -fgpu-rdc. I think we do reduce visibility of non-kernels in that case, but it would be good to make sure.

Checked if we do reduce visibility in such cases, and looks like we do not.

I was indeed mistaken. AFAICT, we only internalize some device-side variables.
@yaxunl - Sam, do we change visibility for anything else? I know we must keep kernels visible, but the question is whether we ever internalize non-kernel functions and if not, whether we want to. In this case it would allow us to bump argument and return value alignment.

For HIP, we mark non-kernel device functions with hidden visibility and internalize them in a LLVM pass for -fno-gpu-rdc.

[tra]

In D120129#3392681, @yaxunl wrote:

For HIP, we mark non-kernel device functions with hidden visibility and internalize them in a LLVM pass for -fno-gpu-rdc.

Looks like now we may have a reason to do so for CUDA, too. Could you point me to where we do it for HIP?

[yaxunl]

In D120129#3392878, @tra wrote:

In D120129#3392681, @yaxunl wrote:

For HIP, we mark non-kernel device functions with hidden visibility and internalize them in a LLVM pass for -fno-gpu-rdc.

Looks like now we may have a reason to do so for CUDA, too. Could you point me to where we do it for HIP?

Make default visibility to be hidden: https://github.com/llvm/llvm-project/blob/main/clang/lib/Driver/ToolChains/HIPAMD.cpp#L203

To avoid making kernels invisible, make them protected visibility: https://github.com/llvm/llvm-project/blob/main/clang/lib/CodeGen/TargetInfo.cpp#L9315

Tell the backend that it needs to internalize non-kernel functions: https://github.com/llvm/llvm-project/blob/main/clang/lib/Driver/ToolChains/HIPAMD.cpp#L189

Let backend internalize non-kernel functions but not variables: https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp#L702

[tra]

In D120129#3396557, @yaxunl wrote:

Thank you for the details!

[kovdan01]

@tra Added appropriate tests for kernels. Looks like that kernel functions get external linkage regardless static specifier, so their parameters alignment is kept untouched.
To ensure that it's true, altered clang/test/CodeGenCUDA/device-fun-linkage.cu and added appropriate assertion in getFunctionParamOptimizedAlign - it checks module metadata and identifies if the function is marked as kernel or not.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

4322

Sure, I meant that :)

llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp

235

I agree, your variant sound better. Fixed.

[tra]

Few cosmetic&style nits. Otherwise LGTM.

clang/test/CodeGenCUDA/device-fun-linkage.cu

23–24 ↗

(On Diff #417265)

Negative testing is tricky. The check here only ensures that the function does not appear between the positive matches around them. If we emit this function as the first or last function in the file, these checks will succeed, even though they should not.

One way to deal with this is to run negative tests separately, with their own RUN lines and check tags.

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

1421

I'd inline it into printout, too.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

1307

This could be moved out of the loop.

1321

No need for the variable, just inline the call into printout statement. Makes it easier to see what we're printing.

1362

Align ParamByValAlign = std::max(Outs[OIdx].Flags.getNonZeroByValAlign(), 
                                                           getFunctionParamOptimizedAlign(F, ETy, DL));

We could try inlining it into the printout, but it may be too unwieldy for that. Up to you.

1364

ditto.

[kovdan01]

clang/test/CodeGenCUDA/device-fun-linkage.cu

23–24 ↗

(On Diff #417265)

Used the approach with separate RUN lines, thanks!

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

1362

Looks too unwieldy for inlining IMHO, so kept in current state.

[jdoerfert]

Our internal CI flagged an assertion in llvm::NVPTXTargetLowering::getFunctionParamOptimizedAlign(llvm::Function const*, llvm::Type*, llvm::DataLayout const&) last night.

Given the error:
static bool llvm::isa_impl_cl<llvm::ConstantAsMetadata, const llvm::Metadata *>::doit(const From *) [To = llvm::ConstantAsMetadata, From = const llvm::Metadata *]: Assertion Val && "isa<> used on a null pointer"' failed.`
I assume it's this patch and the constant as metadata casts.

I'm working on the reproducer now.

[jdoerfert]

Godbold seems to not have included this:

$ cat test.cpp
int main(){return 0;}
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target -march=sm_80 test.cpp
... static bool llvm::isa_impl_cl<llvm::ConstantAsMetadata, const llvm::Metadata *>::doit(const From *) [To = llvm::ConstantAsMetadata, From = const llvm::Metadata *]: Assertion `Val && "isa<> used on a null pointer"' failed.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

4333

FWIW, if (!F->hasLocalLinkage()) { return Align(ABITypeAlign); } <case code>.

[kovdan01]

In D120129#3408088, @jdoerfert wrote:

Godbold seems to not have included this:

$ cat test.cpp
int main(){return 0;}
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target -march=sm_80 test.cpp
... static bool llvm::isa_impl_cl<llvm::ConstantAsMetadata, const llvm::Metadata *>::doit(const From *) [To = llvm::ConstantAsMetadata, From = const llvm::Metadata *]: Assertion `Val && "isa<> used on a null pointer"' failed.

Thanks for the reproducer! Is the problem specific for sm_80 or not? I am unable to reproduce the issue with sm_75 and lower.

[jdoerfert]

In D120129#3408116, @kovdan01 wrote:

In D120129#3408088, @jdoerfert wrote:

Godbold seems to not have included this:

$ cat test.cpp
int main(){return 0;}
$ clang++ -fopenmp -fopenmp-targets=nvptx64 -Xopenmp-target -march=sm_80 test.cpp
... static bool llvm::isa_impl_cl<llvm::ConstantAsMetadata, const llvm::Metadata *>::doit(const From *) [To = llvm::ConstantAsMetadata, From = const llvm::Metadata *]: Assertion `Val && "isa<> used on a null pointer"' failed.

Thanks for the reproducer! Is the problem specific for sm_80 or not? I am unable to reproduce the issue with sm_75 and lower.

You don't need to run it. If you use this command line it doesn't crash?

[kovdan01]

In D120129#3408168, @jdoerfert wrote:

You don't need to run it. If you use this command line it doesn't crash?

Yes, I run the same command (with sm_75 instead of sm_80 because I don't have libomptarget-nvptx-sm_80.bc), and it doesn't crash.

[jdoerfert]

In D120129#3408174, @kovdan01 wrote:

In D120129#3408168, @jdoerfert wrote:

You don't need to run it. If you use this command line it doesn't crash?

Yes, I run the same command (with sm_75 instead of sm_80 because I don't have libomptarget-nvptx-sm_80.bc), and it doesn't crash.

So it's not sm_80 but the cuda version that is important. 11.4.0 works fine for me, 11.0.2 breaks also for sm_70.

[jdoerfert]

Here are the null pointers that cause the assertion:

!nvvm.annotations = !{!8, !9, !8, !10, !10, !10, !10, !11, !11, !10}

!8 = !{null, !"align", i32 8}
!9 = !{null, !"align", i32 8, !"align", i32 65544, !"align", i32 131080}
!10 = !{null, !"align", i32 16}
!11 = !{null, !"align", i32 16, !"align", i32 65552, !"align", i32 131088}

[jdoerfert]

Please revert this commit. It breaks any code using the cuda.11.0.2 libdevice.bc file, the source of the null nvvm.annotations and the annotations with 5 arguments. See below.
Once a fix is in, and a test is added, we can reapply this.

$ opt -S .../cuda-11.0.2/nvvm/libdevice/libdevice.10.bc | tail
!nvvm.annotations = !{!1, !2, !1, !3, !3, !3, !3, !4, !4, !3}

!0 = !{i32 1, i32 4}
!1 = !{null, !"align", i32 8}
!2 = !{null, !"align", i32 8, !"align", i32 65544, !"align", i32 131080}
!3 = !{null, !"align", i32 16}
!4 = !{null, !"align", i32 16, !"align", i32 65552, !"align", i32 131088}
!5 = distinct !{!5, !6}

[kovdan01]

In D120129#3410479, @jdoerfert wrote:

Please revert this commit. It breaks any code using the cuda.11.0.2 libdevice.bc file, the source of the null nvvm.annotations and the annotations with 5 arguments. See below.
Once a fix is in, and a test is added, we can reapply this.

$ opt -S .../cuda-11.0.2/nvvm/libdevice/libdevice.10.bc | tail
!nvvm.annotations = !{!1, !2, !1, !3, !3, !3, !3, !4, !4, !3}

!0 = !{i32 1, i32 4}
!1 = !{null, !"align", i32 8}
!2 = !{null, !"align", i32 8, !"align", i32 65544, !"align", i32 131080}
!3 = !{null, !"align", i32 16}
!4 = !{null, !"align", i32 16, !"align", i32 65552, !"align", i32 131088}
!5 = distinct !{!5, !6}

Got it, will revert the patch shortly. Could you please describe what do you mean by "break" in a bit more detail? Is that only about poorly designed and failing assertions in getFunctionParamOptimizedAlign or there are some other (maybe even functional) issues? If that's only about assertions, won't using isKernelFunction (as suggested in D122550) help?

[jdoerfert]

In D120129#3410493, @kovdan01 wrote:

In D120129#3410479, @jdoerfert wrote:

Please revert this commit. It breaks any code using the cuda.11.0.2 libdevice.bc file, the source of the null nvvm.annotations and the annotations with 5 arguments. See below.
Once a fix is in, and a test is added, we can reapply this.

$ opt -S .../cuda-11.0.2/nvvm/libdevice/libdevice.10.bc | tail
!nvvm.annotations = !{!1, !2, !1, !3, !3, !3, !3, !4, !4, !3}

!0 = !{i32 1, i32 4}
!1 = !{null, !"align", i32 8}
!2 = !{null, !"align", i32 8, !"align", i32 65544, !"align", i32 131080}
!3 = !{null, !"align", i32 16}
!4 = !{null, !"align", i32 16, !"align", i32 65552, !"align", i32 131088}
!5 = distinct !{!5, !6}

Got it, will revert the patch shortly. Could you please describe what do you mean by "break" in a bit more detail? Is that only about poorly designed and failing assertions in getFunctionParamOptimizedAlign or there are some other (maybe even functional) issues? If that's only about assertions, won't using isKernelFunction (as suggested in D122550) help?

The two assertions introduced here do not hold for the libdevice.bc above. So whenever we link the above we will cause the assertions to fail. That's what I mean with break. isKernelFunction is probably going to fix all that. For a test, copy annotations like the above into one of the .ll files and make sure it contains a private/internal function w/ arguments as well.

[kovdan01]

In D120129#3410510, @jdoerfert wrote:

The two assertions introduced here do not hold for the libdevice.bc above. So whenever we link the above we will cause the assertions to fail. That's what I mean with break. isKernelFunction is probably going to fix all that. For a test, copy annotations like the above into one of the .ll files and make sure it contains a private/internal function w/ arguments as well.

OK, thanks for the explanation! Can we just submit a new patch with a fix (like D122550) instead of reverting this one? The problem with revert is that we should also revert D122381 which depends on this patch. Also, adding tests for null nvvm.annotations and the annotations with 5 arguments IMHO will look better when submitted as a separate patch. So, is revert crucial for you or can we just submit a fix separately?

[jdoerfert]

In D120129#3410512, @kovdan01 wrote:

In D120129#3410510, @jdoerfert wrote:

The two assertions introduced here do not hold for the libdevice.bc above. So whenever we link the above we will cause the assertions to fail. That's what I mean with break. isKernelFunction is probably going to fix all that. For a test, copy annotations like the above into one of the .ll files and make sure it contains a private/internal function w/ arguments as well.

OK, thanks for the explanation! Can we just submit a new patch with a fix (like D122550) instead of reverting this one? The problem with revert is that we should also revert D122381 which depends on this patch. Also, adding tests for null nvvm.annotations and the annotations with 5 arguments IMHO will look better when submitted as a separate patch. So, is revert crucial for you or can we just submit a fix separately?

Our internal build bots and CI for some projects are broken for 3 days. I wish to unbreak them so we get actual meaningful results, e.g., see if something else is breaking our build. I'm fine with a separate patch on top but I would prefer it now so people can run code on Monday.

[kovdan01]

In D120129#3410514, @jdoerfert wrote:

Our internal build bots and CI for some projects are broken for 3 days. I wish to unbreak them so we get actual meaningful results, e.g., see if something else is breaking our build. I'm fine with a separate patch on top but I would prefer it now so people can run code on Monday.

OK, I'll submit a separate patch within an hour, and if it looks good to you - will merge it immediately. The fix looks simple enough so there is no need to wait for additional approvals.

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
   Expr = MCSymbolRefExpr::create(Symbol, MCSymbolRefExpr::VK_None,
                                  OutContext);
   return MCOperand::createExpr(Expr);
 }
 
 void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
   const DataLayout &DL = getDataLayout();
   const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
-  const TargetLowering *TLI = STI.getTargetLowering();
+  const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
 
   Type *Ty = F->getReturnType();
 
   bool isABI = (STI.getSmVersion() >= 20);
 
   if (Ty->getTypeID() == Type::VoidTyID)
     return;
 
if (Ty->isFloatingPointTy() || (Ty->isIntegerTy() && !Ty->isIntegerTy(128))) {
       O << ".param .b" << size << " func_retval0";
     } else if (isa<PointerType>(Ty)) {
       O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
         << " func_retval0";
     } else if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
       unsigned totalsz = DL.getTypeAllocSize(Ty);
       unsigned retAlignment = 0;
       if (!getAlign(*F, 0, retAlignment))
-        retAlignment = DL.getABITypeAlignment(Ty);
+        retAlignment = TLI->getFunctionParamOptimizedAlign(F, Ty, DL).value();
       O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
         << "]";
     } else
       llvm_unreachable("Unknown return type");
   } else {
     SmallVector<EVT, 16> vtparts;
     ComputeValueVTs(*TLI, DL, Ty, vtparts);
     unsigned idx = 0;
void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
   getSymbol(I->getParent())->print(O, MAI);
   O << "_param_" << paramIndex;
 }
 
 void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
   const DataLayout &DL = getDataLayout();
   const AttributeList &PAL = F->getAttributes();
   const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
-  const TargetLowering *TLI = STI.getTargetLowering();
+  const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
+
   Function::const_arg_iterator I, E;
   unsigned paramIndex = 0;
   bool first = true;
   bool isKernelFunc = isKernelFunction(*F);
   bool isABI = (STI.getSmVersion() >= 20);
   bool hasImageHandles = STI.hasImageHandles();
   MVT thePointerTy = TLI->getPointerTy(DL);
 
if (isKernelFunction(*F)) {
             O << "\t.param .samplerref ";
           CurrentFnSym->print(O, MAI);
           O << "_param_" << paramIndex;
         }
         continue;
       }
     }
 
+    auto getOptimalAlignForParam = [TLI, &DL, &PAL, F,
+                                    paramIndex](Type *Ty) -> Align {
+      Align TypeAlign = TLI->getFunctionParamOptimizedAlign(F, Ty, DL);
+      MaybeAlign ParamAlign = PAL.getParamAlignment(paramIndex);
+      return max(TypeAlign, ParamAlign);
+    };
+
     if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
       if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
         // Just print .param .align <a> .b8 .param[size];
-        // <a> = PAL.getparamalignment
+        // <a>  = optimal alignment for the element type; always multiple of
+        //        PAL.getParamAlignment
         // size = typeallocsize of element type
-        const Align align = DL.getValueOrABITypeAlignment(
-            PAL.getParamAlignment(paramIndex), Ty);
+        Align OptimalAlign = getOptimalAlignForParam(Ty);
 
         unsigned sz = DL.getTypeAllocSize(Ty);
-        O << "\t.param .align " << align.value() << " .b8 ";
+        O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
         printParamName(I, paramIndex, O);
         O << "[" << sz << "]";
 
         continue;
       }
       // Just a scalar
       auto *PTy = dyn_cast<PointerType>(Ty);
       if (isKernelFunc) {
for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
     }
 
     // param has byVal attribute.
     Type *ETy = PAL.getParamByValType(paramIndex);
     assert(ETy && "Param should have byval type");
 
     if (isABI || isKernelFunc) {
       // Just print .param .align <a> .b8 .param[size];
-      // <a> = PAL.getparamalignment
+      // <a>  = optimal alignment for the element type; always multiple of
+      //        PAL.getParamAlignment
       // size = typeallocsize of element type
-      Align align =
-          DL.getValueOrABITypeAlignment(PAL.getParamAlignment(paramIndex), ETy);
+      Align OptimalAlign = getOptimalAlignForParam(ETy);
+
       // Work around a bug in ptxas. When PTX code takes address of
       // byval parameter with alignment < 4, ptxas generates code to
       // spill argument into memory. Alas on sm_50+ ptxas generates
       // SASS code that fails with misaligned access. To work around
       // the problem, make sure that we align byval parameters by at
       // least 4. Matching change must be made in LowerCall() where we
       // prepare parameters for the call.
       //
       // TODO: this will need to be undone when we get to support multi-TU
       // device-side compilation as it breaks ABI compatibility with nvcc.
       // Hopefully ptxas bug is fixed by then.
-      if (!isKernelFunc && align < Align(4))
-        align = Align(4);
+      if (!isKernelFunc && OptimalAlign < Align(4))
+        OptimalAlign = Align(4);
       unsigned sz = DL.getTypeAllocSize(ETy);
-      O << "\t.param .align " << align.value() << " .b8 ";
+      O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
       printParamName(I, paramIndex, O);
       O << "[" << sz << "]";
       continue;
     } else {
       // Split the ETy into constituent parts and
       // print .param .b<size> <name> for each part.
       // Further, if a part is vector, print the above for
       // each vector element.

llvm/lib/Target/NVPTX/NVPTXISelLowering.h

public:
   SDValue LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const;
 
   const char *getTargetNodeName(unsigned Opcode) const override;
 
   bool getTgtMemIntrinsic(IntrinsicInfo &Info, const CallInst &I,
                           MachineFunction &MF,
                           unsigned Intrinsic) const override;
 
+  /// getFunctionParamOptimizedAlign - since function arguments are passed via
+  /// .param space, we may want to have increased alignmet in it to enhance
+  /// vectorization options. We can increase alignment only if the function has
+  /// internal or private linkage as for other linkage types callers may already
+  /// rely on default linkage. This function returns default alignment
+  /// multiplied by 4, 2 or 1, choosing the maximum multiplier that alignment
+  /// remains type size divisor.
+  Align getFunctionParamOptimizedAlign(const Function *F, Type *ArgTy,
+                                       const DataLayout &DL) const;
+
   /// isLegalAddressingMode - Return true if the addressing mode represented
   /// by AM is legal for this target, for a load/store of the specified type
   /// Used to guide target specific optimizations, like loop strength
   /// reduction (LoopStrengthReduce.cpp) and memory optimization for
   /// address mode (CodeGenPrepare.cpp)
   bool isLegalAddressingMode(const DataLayout &DL, const AddrMode &AM, Type *Ty,
                              unsigned AS,
                              Instruction *I = nullptr) const override;

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

if (retTy->getTypeID() == Type::VoidTyID) {
     O << ") ";
   }
   O << "_ (";
 
   bool first = true;
 
   unsigned OIdx = 0;
   for (unsigned i = 0, e = Args.size(); i != e; ++i, ++OIdx) {
+    const Function *F = CB.getFunction();
     Type *Ty = Args[i].Ty;
     if (!first) {
       O << ", ";
     }
     first = false;
 
     if (!Outs[OIdx].Flags.isByVal()) {
       if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
-        unsigned align = 0;
+        unsigned ParamAlign = 0;
         const CallInst *CallI = cast<CallInst>(&CB);
         // +1 because index 0 is reserved for return type alignment
-        if (!getAlign(*CallI, i + 1, align))
-          align = DL.getABITypeAlignment(Ty);
+        if (!getAlign(*CallI, i + 1, ParamAlign))
+          ParamAlign = getFunctionParamOptimizedAlign(F, Ty, DL).value();
         unsigned sz = DL.getTypeAllocSize(Ty);
-        O << ".param .align " << align << " .b8 ";
+        O << ".param .align " << ParamAlign << " .b8 ";
         O << "_";
         O << "[" << sz << "]";
         // update the index for Outs
         SmallVector<EVT, 16> vtparts;
         ComputeValueVTs(*this, DL, Ty, vtparts);
         if (unsigned len = vtparts.size())
           OIdx += len - 1;
         continue;
if (!Outs[OIdx].Flags.isByVal()) {
         sz = 32;
       else
         sz = Ty->getPrimitiveSizeInBits();
       O << ".param .b" << sz << " ";
       O << "_";
       continue;
     }
 
-    Align align = Outs[OIdx].Flags.getNonZeroByValAlign();
+    Align ParamByValAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
+
+    // Try to increase alignment. This code matches logic in LowerCall when
+    // alignment increase is performed to increase vectorization options.
+    Type *ETy = Args[i].IndirectType;
+    Align AlignCandidate = getFunctionParamOptimizedAlign(F, ETy, DL);
+    ParamByValAlign = std::max(ParamByValAlign, AlignCandidate);
+
     unsigned sz = Outs[OIdx].Flags.getByValSize();
-    O << ".param .align " << align.value() << " .b8 ";
+    O << ".param .align " << ParamByValAlign.value() << " .b8 ";
     O << "_";
     O << "[" << sz << "]";
   }
   O << ");";
   return O.str();
 }
 
 Align NVPTXTargetLowering::getArgumentAlignment(SDValue Callee,
if (const auto *CI = dyn_cast<CallInst>(CB)) {
       // Function?
       if (const auto *CalleeF = dyn_cast<Function>(CalleeV))
         DirectCallee = CalleeF;
     }
   }
 
   // Check for function alignment information if we found that the
   // ultimate target is a Function
-  if (DirectCallee)
+  if (DirectCallee) {
     if (getAlign(*DirectCallee, Idx, Alignment))
       return Align(Alignment);
+    // If alignment information is not available, fall back to the
+    // default function param optimized type alignment
+    return getFunctionParamOptimizedAlign(DirectCallee, Ty, DL);
+  }
 
-  // Call is indirect or alignment information is not available, fall back to
-  // the ABI type alignment
+  // Call is indirect, fall back to the ABI type alignment
   return DL.getABITypeAlign(Ty);
 }
 
 SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
                                        SmallVectorImpl<SDValue> &InVals) const {
   SelectionDAG &DAG = CLI.DAG;
   SDLoc dl = CLI.DL;
   SmallVectorImpl<ISD::OutputArg> &Outs = CLI.Outs;
SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
   //   * if there is a vector argument with more than typical vector-length
   //     elements (generally if more than 4) where each vector element is
   //     individually present in Outs.
   // So a different index should be used for indexing into Outs/OutVals.
   // See similar issue in LowerFormalArguments.
   unsigned OIdx = 0;
   // Declare the .params or .reg need to pass values
   // to the function
+
   for (unsigned i = 0, e = Args.size(); i != e; ++i, ++OIdx) {
     EVT VT = Outs[OIdx].VT;
     Type *Ty = Args[i].Ty;
 
     if (!Outs[OIdx].Flags.isByVal()) {
       SmallVector<EVT, 16> VTs;
       SmallVector<uint64_t, 16> Offsets;
       ComputePTXValueVTs(*this, DL, Ty, VTs, &Offsets);
if (!Outs[OIdx].Flags.isByVal()) {
       assert(StoreOperands.empty() && "Unfinished parameter store.");
       if (VTs.size() > 0)
         --OIdx;
       ++paramCount;
       continue;
     }
 
     // ByVal arguments
+    // TODO: remove code duplication when handling byval and non-byval cases.
     SmallVector<EVT, 16> VTs;
     SmallVector<uint64_t, 16> Offsets;
-    assert(Args[i].IndirectType && "byval arg must have indirect type");
-    ComputePTXValueVTs(*this, DL, Args[i].IndirectType, VTs, &Offsets, 0);
+    Type *ETy = Args[i].IndirectType;
+    assert(ETy && "byval arg must have indirect type");
+    ComputePTXValueVTs(*this, DL, ETy, VTs, &Offsets, 0);
 
     // declare .param .align <align> .b8 .param<n>[<size>];
     unsigned sz = Outs[OIdx].Flags.getByValSize();
     SDVTList DeclareParamVTs = DAG.getVTList(MVT::Other, MVT::Glue);
-    Align ArgAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
+
     // The ByValAlign in the Outs[OIdx].Flags is alway set at this point,
     // so we don't need to worry about natural alignment or not.
     // See TargetLowering::LowerCallTo().
+    Align ArgAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
+
+    // Try to increase alignment to enhance vectorization options.
+    const Function *F = CB->getCalledFunction();
+    Align AlignCandidate = getFunctionParamOptimizedAlign(F, ETy, DL);
+    ArgAlign = std::max(ArgAlign, AlignCandidate);
 
     // Enforce minumum alignment of 4 to work around ptxas miscompile
     // for sm_50+. See corresponding alignment adjustment in
     // emitFunctionParamList() for details.
     if (ArgAlign < Align(4))
       ArgAlign = Align(4);
     SDValue DeclareParamOps[] = {
         Chain, DAG.getConstant(ArgAlign.value(), dl, MVT::i32),
         DAG.getConstant(paramCount, dl, MVT::i32),
         DAG.getConstant(sz, dl, MVT::i32), InFlag};
     Chain = DAG.getNode(NVPTXISD::DeclareParam, dl, DeclareParamVTs,
                         DeclareParamOps);
     InFlag = Chain.getValue(1);
+
+    auto VectorInfo = VectorizePTXValueVTs(VTs, Offsets, ArgAlign);
+    SmallVector<SDValue, 6> StoreOperands;
     for (unsigned j = 0, je = VTs.size(); j != je; ++j) {
       EVT elemtype = VTs[j];
       int curOffset = Offsets[j];
-      unsigned PartAlign = GreatestCommonDivisor64(ArgAlign.value(), curOffset);
+      Align PartAlign = commonAlignment(ArgAlign, curOffset);
+
+      // New store.
+      if (VectorInfo[j] & PVF_FIRST) {
+        assert(StoreOperands.empty() && "Unfinished preceding store.");
+        StoreOperands.push_back(Chain);
+        StoreOperands.push_back(DAG.getConstant(paramCount, dl, MVT::i32));
+        StoreOperands.push_back(DAG.getConstant(curOffset, dl, MVT::i32));
+      }
+
       auto PtrVT = getPointerTy(DL);
       SDValue srcAddr = DAG.getNode(ISD::ADD, dl, PtrVT, OutVals[OIdx],
                                     DAG.getConstant(curOffset, dl, PtrVT));
       SDValue theVal = DAG.getLoad(elemtype, dl, tempChain, srcAddr,
                                    MachinePointerInfo(), PartAlign);
+
       if (elemtype.getSizeInBits() < 16) {
+        // Use 16-bit registers for small stores as it's the
+        // smallest general purpose register size supported by NVPTX.
         theVal = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i16, theVal);
       }
-      SDVTList CopyParamVTs = DAG.getVTList(MVT::Other, MVT::Glue);
-      SDValue CopyParamOps[] = { Chain,
-                                 DAG.getConstant(paramCount, dl, MVT::i32),
-                                 DAG.getConstant(curOffset, dl, MVT::i32),
-                                 theVal, InFlag };
-      Chain = DAG.getMemIntrinsicNode(
-          NVPTXISD::StoreParam, dl, CopyParamVTs, CopyParamOps, elemtype,
-          MachinePointerInfo(), /* Align */ None, MachineMemOperand::MOStore);
 
+      // Record the value to store.
+      StoreOperands.push_back(theVal);
+
+      if (VectorInfo[j] & PVF_LAST) {
+        unsigned NumElts = StoreOperands.size() - 3;
+        NVPTXISD::NodeType Op;
+        switch (NumElts) {
+        case 1:
+          Op = NVPTXISD::StoreParam;
+          break;
+        case 2:
+          Op = NVPTXISD::StoreParamV2;
+          break;
+        case 4:
+          Op = NVPTXISD::StoreParamV4;
+          break;
+        default:
+          llvm_unreachable("Invalid vector info.");
+        }
+
+        StoreOperands.push_back(InFlag);
+
+        Chain = DAG.getMemIntrinsicNode(
+            Op, dl, DAG.getVTList(MVT::Other, MVT::Glue), StoreOperands,
+            elemtype, MachinePointerInfo(), PartAlign,
+            MachineMemOperand::MOStore);
         InFlag = Chain.getValue(1);
+
+        // Cleanup.
+        StoreOperands.clear();
       }
+    }
+    assert(StoreOperands.empty() && "Unfinished parameter store.");
+    if (VTs.size() > 0)
+      --OIdx;
     ++paramCount;
   }
 
   GlobalAddressSDNode *Func = dyn_cast<GlobalAddressSDNode>(Callee.getNode());
   MaybeAlign retAlignment = None;
 
   // Handle Result
   if (Ins.size() > 0) {

 }
 
 SDValue
 NVPTXTargetLowering::LowerReturn(SDValue Chain, CallingConv::ID CallConv,
                                  bool isVarArg,
                                  const SmallVectorImpl<ISD::OutputArg> &Outs,
                                  const SmallVectorImpl<SDValue> &OutVals,
                                  const SDLoc &dl, SelectionDAG &DAG) const {
-  MachineFunction &MF = DAG.getMachineFunction();
+  const MachineFunction &MF = DAG.getMachineFunction();
+  const Function &F = MF.getFunction();
   Type *RetTy = MF.getFunction().getReturnType();
 
   bool isABI = (STI.getSmVersion() >= 20);
   assert(isABI && "Non-ABI compilation is not supported");
   if (!isABI)
     return Chain;
 
   const DataLayout &DL = DAG.getDataLayout();
   SmallVector<EVT, 16> VTs;
   SmallVector<uint64_t, 16> Offsets;
   ComputePTXValueVTs(*this, DL, RetTy, VTs, &Offsets);
   assert(VTs.size() == OutVals.size() && "Bad return value decomposition");
 
   auto VectorInfo = VectorizePTXValueVTs(
-      VTs, Offsets, RetTy->isSized() ? DL.getABITypeAlign(RetTy) : Align(1));
+      VTs, Offsets,
+      RetTy->isSized() ? getFunctionParamOptimizedAlign(&F, RetTy, DL)
+                       : Align(1));
 
   // PTX Interoperability Guide 3.3(A): [Integer] Values shorter than
   // 32-bits are sign extended or zero extended, depending on whether
   // they are signed or unsigned types.
   bool ExtendIntegerRetVal =
       RetTy->isIntegerTy() && DL.getTypeAllocSizeInBits(RetTy) < 32;
 
   SmallVector<SDValue, 6> StoreOperands;
case Intrinsic::nvvm_suld_3d_v2i64_zero:
     Info.offset = 0;
     Info.flags = MachineMemOperand::MOLoad;
     Info.align = Align(16);
     return true;
   }
   return false;
 }
 
+/// getFunctionParamOptimizedAlign - since function arguments are passed via
+/// .param space, we may want to have increased alignmet in it to enhance
+/// vectorization options. We can increase alignment only if the function has
+/// internal or private linkage as for other linkage types callers may already
+/// rely on default linkage. This function returns default alignment
+/// multiplied by 4, 2 or 1, choosing the maximum multiplier that alignment
+/// remains type size divisor.
+Align NVPTXTargetLowering::getFunctionParamOptimizedAlign(
+    const Function *F, Type *ArgTy, const DataLayout &DL) const {
+  const unsigned ABITypeAlign = DL.getABITypeAlign(ArgTy).value();
+
+  // If a function has linkage different from internal or private, we
+  // must use default ABI alignment as external users rely on it.
+  switch (F->getLinkage()) {
+  case GlobalValue::InternalLinkage:
+  case GlobalValue::PrivateLinkage: {
+    unsigned Size = DL.getTypeAllocSize(ArgTy);
+    if (Size % (ABITypeAlign * 4) == 0)
+      return Align(ABITypeAlign * 4);
+    if (Size % (ABITypeAlign * 2) == 0)
+      return Align(ABITypeAlign * 2);
+    return Align(ABITypeAlign);
+  }
+  default:
+    return Align(ABITypeAlign);
+  }
+}
+
 /// isLegalAddressingMode - Return true if the addressing mode represented
 /// by AM is legal for this target, for a load/store of the specified type.
 /// Used to guide target specific optimizations, like loop strength reduction
 /// (LoopStrengthReduce.cpp) and memory optimization for address mode
 /// (CodeGenPrepare.cpp)
 bool NVPTXTargetLowering::isLegalAddressingMode(const DataLayout &DL,
                                                 const AddrMode &AM, Type *Ty,
                                                 unsigned AS, Instruction *I) const {

llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp

 //      %b_generic = addrspacecast i32 addrspace(1)* %b_global to i32*
 //      ; use %b_generic
 //    }
 //
 // TODO: merge this pass with NVPTXInferAddressSpaces so that other passes don't
 // cancel the addrspacecast pair this pass emits.
 //===----------------------------------------------------------------------===//
 
+#include "MCTargetDesc/NVPTXBaseInfo.h"
 #include "NVPTX.h"
 #include "NVPTXTargetMachine.h"
 #include "NVPTXUtilities.h"
-#include "MCTargetDesc/NVPTXBaseInfo.h"
 #include "llvm/Analysis/ValueTracking.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Type.h"
 #include "llvm/Pass.h"
+#include <queue>
 
 #define DEBUG_TYPE "nvptx-lower-args"
 
 using namespace llvm;
 
 namespace llvm {
 void initializeNVPTXLowerArgsPass(PassRegistry &);
 }
static void convertToParamAS(Value *OldUser, Value *Param) {
   // closest to the load are deleted first. Otherwise they may still be in use.
   // E.g if we have Value = Load(BitCast(GEP(arg))), InstructionsToDelete will
   // have {GEP,BitCast}. GEP can't be deleted first, because it's still used by
   // the BitCast.
   llvm::for_each(reverse(InstructionsToDelete),
                  [](Instruction *I) { I->eraseFromParent(); });
 }
 
+// Try to increase alignment in param address space. We can force a
+// specific alignment for byval params to make vectorization possible,
+// e.g. ld.param.v2* and ld.param.v4* instead of ld.param.
+// Such changes must be conformed with parameter store and load in
+// NVPTXTargetLowering::LowerCall.
+static void tryIncreaseByValAlign(Argument *Arg, Value *ArgInParamAS,
+                                  const NVPTXTargetLowering *TLI) {
+  Function *Func = Arg->getParent();
+  Type *StructType = Arg->getParamByValType();
+  const DataLayout DL(Func->getParent());
+
+  uint64_t NewArgAlign =
+      TLI->getFunctionParamOptimizedAlign(Func, StructType, DL).value();
+  uint64_t CurArgAlign =
+      Arg->getAttribute(Attribute::Alignment).getValueAsInt();
+
+  if (CurArgAlign >= NewArgAlign)
+    return;
+
+  LLVM_DEBUG(dbgs() << "Try to use alignment " << NewArgAlign << " instead of "
+                    << CurArgAlign << " for " << *Arg << '\n');
+
+  auto NewAlignAttr =
+      Attribute::get(Func->getContext(), Attribute::Alignment, NewArgAlign);
+  Arg->removeAttr(Attribute::Alignment);
+  Arg->addAttr(NewAlignAttr);
+
+  struct Load {
+    LoadInst *Inst;
+    uint64_t Offset;
+  };
+
+  struct LoadContext {
+    Value *InitialVal;
+    uint64_t Offset;
+  };
+
+  SmallVector<Load> Loads;
+  std::queue<LoadContext> Worklist;
+  Worklist.push({ArgInParamAS, 0});
+
+  while (!Worklist.empty()) {
+    LoadContext Ctx = Worklist.front();
+    Worklist.pop();
+
+    for (User *CurUser : Ctx.InitialVal->users()) {
+      if (auto *I = dyn_cast<LoadInst>(CurUser)) {
+        Loads.push_back({I, Ctx.Offset});
+        continue;
+      }
+
+      if (auto *I = dyn_cast<BitCastInst>(CurUser)) {
+        Worklist.push({I, Ctx.Offset});
+        continue;
+      }
+
+      if (auto *I = dyn_cast<GetElementPtrInst>(CurUser)) {
+        APInt OffsetAccumulated =
+            APInt::getZero(DL.getIndexSizeInBits(ADDRESS_SPACE_PARAM));
+
+        if (!I->accumulateConstantOffset(DL, OffsetAccumulated))
+          continue;
+
+        uint64_t OffsetLimit = -1;
+        uint64_t Offset = OffsetAccumulated.getLimitedValue(OffsetLimit);
+        assert(Offset != OffsetLimit && "Expect Offset less than UINT64_MAX");
+
+        Worklist.push({I, Ctx.Offset + Offset});
+        continue;
+      }
+
+      llvm_unreachable("All users must be one of: load, "
+                       "bitcast, getelementptr.");
+    }
+  }
+
+  for (Load &CurLoad : Loads) {
+    Align NewLoadAlign(greatestCommonDivisor(NewArgAlign, CurLoad.Offset));
+    Align CurLoadAlign(CurLoad.Inst->getAlign());
+    CurLoad.Inst->setAlignment(std::max(NewLoadAlign, CurLoadAlign));
+  }
+}
+
 void NVPTXLowerArgs::handleByValParam(Argument *Arg) {
   Function *Func = Arg->getParent();
   Instruction *FirstInst = &(Func->getEntryBlock().front());
   Type *StructType = Arg->getParamByValType();
   assert(StructType && "Missing byval type");
 
   auto IsALoadChain = [&](Value *Start) {
     SmallVector<Value *, 16> ValuesToCheck = {Start};
if (llvm::all_of(Arg->users(), IsALoadChain)) {
     SmallVector<User *, 16> UsersToUpdate(Arg->users());
     Value *ArgInParamAS = new AddrSpaceCastInst(
         Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
         FirstInst);
     llvm::for_each(UsersToUpdate, [ArgInParamAS](Value *V) {
       convertToParamAS(V, ArgInParamAS);
     });
     LLVM_DEBUG(dbgs() << "No need to copy " << *Arg << "\n");
+
+    // Further optimizations require target lowering info.
+    if (!TM)
+      return;
+
+    const auto *TLI =
+        cast<NVPTXTargetLowering>(TM->getSubtargetImpl()->getTargetLowering());
+
+    tryIncreaseByValAlign(Arg, ArgInParamAS, TLI);
+
     return;
   }
 
   // Otherwise we have to create a temporary copy.
   const DataLayout &DL = Func->getParent()->getDataLayout();
   unsigned AS = DL.getAllocaAddrSpace();
   AllocaInst *AllocA = new AllocaInst(StructType, AS, Arg->getName(), FirstInst);
   // Set the alignment to alignment of the byval parameter. This is because,

llvm/test/CodeGen/NVPTX/param-vectorize.ll

+; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s
+;
+; This is IR generated with clang using -O3 optimization level
+; and nvptx-unknown-unknown target from the following C code.
+;
+; struct St4  { unsigned field[1]; };
+; struct St8  { unsigned field[2]; };
+; struct St12 { unsigned field[3]; };
+; struct St16 { unsigned field[4]; };
+; struct St20 { unsigned field[5]; };
+; struct St24 { unsigned field[6]; };
+; struct St28 { unsigned field[7]; };
+; struct St32 { unsigned field[8]; };
+; 
+; #define DECLARE_CALLEE(StName)                                  \
+; __attribute__((noinline))                                       \
+; struct StName callee_##StName(struct StName in) {               \
+;   struct StName ret;                                            \
+;   const unsigned size = sizeof(ret.field) / sizeof(unsigned);   \
+;   for (unsigned i = 0; i != size; ++i)                          \
+;     ret.field[i] = in.field[i];                                 \
+;   return ret;                                                   \
+; }                                                               \
+; 
+; #define DECLARE_CALLER(StName)                      \
+; struct StName caller_##StName(struct StName* in)    \
+; {                                                   \
+;   struct StName out = callee_##StName(*in);         \
+;   return out;                                       \
+; }                                                   \
+; 
+; #define DECLARE_CALL(StName)    \
+;     DECLARE_CALLEE(StName)      \
+;     DECLARE_CALLER(StName)      \
+; 
+; DECLARE_CALL(St4)
+; DECLARE_CALL(St8)
+; DECLARE_CALL(St12)
+; DECLARE_CALL(St16)
+; DECLARE_CALL(St20)
+; DECLARE_CALL(St24)
+; DECLARE_CALL(St28)
+; DECLARE_CALL(St32)
+
+%struct.St4 = type { [1 x i32] }
+%struct.St8 = type { [2 x i32] }
+%struct.St12 = type { [3 x i32] }
+%struct.St16 = type { [4 x i32] }
+%struct.St20 = type { [5 x i32] }
+%struct.St24 = type { [6 x i32] }
+%struct.St28 = type { [7 x i32] }
+%struct.St32 = type { [8 x i32] }
+
+; Section 1 - checking that:
+; - function argument (including retval) vectorization is done with internal linkage;
+; - caller and callee specify correct alignment for callee's params;
+; - caller with default (external) linkage has st.param unvectorized and its retval has default alignment.
+
+define internal %struct.St4 @callee_St4(%struct.St4* nocapture noundef readonly byval(%struct.St4) align 4 %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[4])
+  ; CHECK-LABEL: callee_St4(
+  ; CHECK-NEXT:  .param .align 4 .b8 callee_St4_param_0[4]
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [callee_St4_param_0];
+  ; CHECK:       st.param.b32    [func_retval0+0], [[R1]];
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St4, %struct.St4* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %.fca.0.0.insert = insertvalue %struct.St4 poison, i32 %1, 0, 0
+  ret %struct.St4 %.fca.0.0.insert
+}
+
+define %struct.St4 @caller_St4(%struct.St4* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[4])
+  ; CHECK-LABEL: caller_St4(
+  ; CHECK-NEXT:  .param .b32 caller_St4_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St4_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       .param .align 4 .b8 param0[4];
+  ; CHECK:       st.param.b32    [param0+0], [[R2]];
+  ; CHECK:       .param .align 4 .b8 retval0[4];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St4, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.b32    [[R3:%r[0-9]+]], [retval0+0];
+  ; CHECK:       st.param.b32    [func_retval0+0], [[R3]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St4 @callee_St4(%struct.St4* noundef byval(%struct.St4) align 4 %in)
+  %1 = extractvalue %struct.St4 %call, 0
+  %.fca.0.extract = extractvalue [1 x i32] %1, 0
+  %.fca.0.0.insert = insertvalue %struct.St4 poison, i32 %.fca.0.extract, 0, 0
+  ret %struct.St4 %.fca.0.0.insert
+}
+
+define internal %struct.St8 @callee_St8(%struct.St8* nocapture noundef readonly byval(%struct.St8) align 4 %in) {
+  ; CHECK:       .func  (.param .align 8 .b8 func_retval0[8])
+  ; CHECK-LABEL: callee_St8(
+  ; CHECK-NEXT:  .param .align 8 .b8 callee_St8_param_0[8]
+  ; CHECK:       ld.param.v2.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St8_param_0];
+  ; CHECK:       st.param.v2.b32 [func_retval0+0], {[[R1]], [[R2]]};
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St8, %struct.St8* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St8, %struct.St8* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %.fca.0.0.insert = insertvalue %struct.St8 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St8 %.fca.0.0.insert, i32 %2, 0, 1
+  ret %struct.St8 %.fca.0.1.insert
+}
+
+define %struct.St8 @caller_St8(%struct.St8* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[8])
+  ; CHECK-LABEL: caller_St8(
+  ; CHECK-NEXT:  .param .b32 caller_St8_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St8_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       .param .align 8 .b8 param0[8];
+  ; CHECK:       st.param.v2.b32 [param0+0], {[[R3]], [[R2]]};
+  ; CHECK:       .param .align 8 .b8 retval0[8];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St8, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.v2.b32 {[[R4:%r[0-9]+]], [[R5:%r[0-9]+]]}, [retval0+0];
+  ; CHECK:       st.param.b32    [func_retval0+0], [[R4]];
+  ; CHECK:       st.param.b32    [func_retval0+4], [[R5]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St8 @callee_St8(%struct.St8* noundef byval(%struct.St8) align 4 %in)
+  %1 = extractvalue %struct.St8 %call, 0
+  %.fca.0.extract = extractvalue [2 x i32] %1, 0
+  %.fca.1.extract = extractvalue [2 x i32] %1, 1
+  %.fca.0.0.insert = insertvalue %struct.St8 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St8 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  ret %struct.St8 %.fca.0.1.insert
+}
+
+define internal %struct.St12 @callee_St12(%struct.St12* nocapture noundef readonly byval(%struct.St12) align 4 %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[12])
+  ; CHECK-LABEL: callee_St12(
+  ; CHECK-NEXT:  .param .align 4 .b8 callee_St12_param_0[12]
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]],  [callee_St12_param_0];
+  ; CHECK:       ld.param.u32    [[R2:%r[0-9]+]],  [callee_St12_param_0+4];
+  ; CHECK:       ld.param.u32    [[R3:%r[0-9]+]],  [callee_St12_param_0+8];
+  ; CHECK:       st.param.b32    [func_retval0+0], [[R1]];
+  ; CHECK:       st.param.b32    [func_retval0+4], [[R2]];
+  ; CHECK:       st.param.b32    [func_retval0+8], [[R3]];
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St12, %struct.St12* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St12, %struct.St12* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St12, %struct.St12* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %.fca.0.0.insert = insertvalue %struct.St12 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St12 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St12 %.fca.0.1.insert, i32 %3, 0, 2
+  ret %struct.St12 %.fca.0.2.insert
+}
+
+define %struct.St12 @caller_St12(%struct.St12* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[12])
+  ; CHECK-LABEL: caller_St12(
+  ; CHECK-NEXT:  .param .b32 caller_St12_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St12_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       .param .align 4 .b8 param0[12];
+  ; CHECK:       st.param.b32    [param0+0], [[R2]];
+  ; CHECK:       st.param.b32    [param0+4], [[R3]];
+  ; CHECK:       st.param.b32    [param0+8], [[R4]];
+  ; CHECK:       .param .align 4 .b8 retval0[12];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St12, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.b32    [[R5:%r[0-9]+]],  [retval0+0];
+  ; CHECK:       ld.param.b32    [[R6:%r[0-9]+]],  [retval0+4];
+  ; CHECK:       ld.param.b32    [[R7:%r[0-9]+]],  [retval0+8];
+  ; CHECK:       st.param.b32    [func_retval0+0], [[R5]];
+  ; CHECK:       st.param.b32    [func_retval0+4], [[R6]];
+  ; CHECK:       st.param.b32    [func_retval0+8], [[R7]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St12 @callee_St12(%struct.St12* noundef byval(%struct.St12) align 4 %in)
+  %1 = extractvalue %struct.St12 %call, 0
+  %.fca.0.extract = extractvalue [3 x i32] %1, 0
+  %.fca.1.extract = extractvalue [3 x i32] %1, 1
+  %.fca.2.extract = extractvalue [3 x i32] %1, 2
+  %.fca.0.0.insert = insertvalue %struct.St12 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St12 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St12 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  ret %struct.St12 %.fca.0.2.insert
+}
+
+define internal %struct.St16 @callee_St16(%struct.St16* nocapture noundef readonly byval(%struct.St16) align 4 %in) {
+  ; CHECK:       .func  (.param .align 16 .b8 func_retval0[16])
+  ; CHECK-LABEL: callee_St16(
+  ; CHECK-NEXT:  .param .align 16 .b8 callee_St16_param_0[16]
+  ; CHECK:       ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St16_param_0];
+  ; CHECK:       st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %.fca.0.0.insert = insertvalue %struct.St16 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St16 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St16 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St16 %.fca.0.2.insert, i32 %4, 0, 3
+  ret %struct.St16 %.fca.0.3.insert
+}
+
+define %struct.St16 @caller_St16(%struct.St16* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[16])
+  ; CHECK-LABEL: caller_St16(
+  ; CHECK-NEXT:  .param .b32 caller_St16_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St16_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]+12];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R5:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       .param .align 16 .b8 param0[16];
+  ; CHECK:       st.param.v4.b32 [param0+0], {[[R5]], [[R4]], [[R3]], [[R2]]};
+  ; CHECK:       .param .align 16 .b8 retval0[16];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St16, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.v4.b32 {[[R6:%r[0-9]+]], [[R7:%r[0-9]+]], [[R8:%r[0-9]+]], [[R9:%r[0-9]+]]}, [retval0+0];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R6]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R7]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R8]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R9]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St16 @callee_St16(%struct.St16* noundef byval(%struct.St16) align 4 %in)
+  %1 = extractvalue %struct.St16 %call, 0
+  %.fca.0.extract = extractvalue [4 x i32] %1, 0
+  %.fca.1.extract = extractvalue [4 x i32] %1, 1
+  %.fca.2.extract = extractvalue [4 x i32] %1, 2
+  %.fca.3.extract = extractvalue [4 x i32] %1, 3
+  %.fca.0.0.insert = insertvalue %struct.St16 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St16 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St16 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St16 %.fca.0.2.insert, i32 %.fca.3.extract, 0, 3
+  ret %struct.St16 %.fca.0.3.insert
+}
+
+define internal %struct.St20 @callee_St20(%struct.St20* nocapture noundef readonly byval(%struct.St20) align 4 %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[20])
+  ; CHECK-LABEL: callee_St20(
+  ; CHECK-NEXT:  .param .align 4 .b8 callee_St20_param_0[20]
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]],   [callee_St20_param_0];
+  ; CHECK:       ld.param.u32    [[R2:%r[0-9]+]],   [callee_St20_param_0+4];
+  ; CHECK:       ld.param.u32    [[R3:%r[0-9]+]],   [callee_St20_param_0+8];
+  ; CHECK:       ld.param.u32    [[R4:%r[0-9]+]],   [callee_St20_param_0+12];
+  ; CHECK:       ld.param.u32    [[R5:%r[0-9]+]],   [callee_St20_param_0+16];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R1]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R2]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R3]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R4]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R5]];
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St20, %struct.St20* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St20, %struct.St20* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St20, %struct.St20* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St20, %struct.St20* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %arrayidx.4 = getelementptr inbounds %struct.St20, %struct.St20* %in, i32 0, i32 0, i32 4
+  %5 = load i32, i32* %arrayidx.4, align 4
+  %.fca.0.0.insert = insertvalue %struct.St20 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St20 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St20 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St20 %.fca.0.2.insert, i32 %4, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St20 %.fca.0.3.insert, i32 %5, 0, 4
+  ret %struct.St20 %.fca.0.4.insert
+}
+
+define %struct.St20 @caller_St20(%struct.St20* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[20])
+  ; CHECK-LABEL: caller_St20(
+  ; CHECK-NEXT:  .param .b32 caller_St20_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St20_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       ld.u32          [[R5:%r[0-9]+]], [[[R1]]+12];
+  ; CHECK:       ld.u32          [[R6:%r[0-9]+]], [[[R1]]+16];
+  ; CHECK:       .param .align 4 .b8 param0[20];
+  ; CHECK:       st.param.b32    [param0+0],  [[R2]];
+  ; CHECK:       st.param.b32    [param0+4],  [[R3]];
+  ; CHECK:       st.param.b32    [param0+8],  [[R4]];
+  ; CHECK:       st.param.b32    [param0+12], [[R5]];
+  ; CHECK:       st.param.b32    [param0+16], [[R6]];
+  ; CHECK:       .param .align 4 .b8 retval0[20];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St20, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.b32    [[R7:%r[0-9]+]],   [retval0+0];
+  ; CHECK:       ld.param.b32    [[R8:%r[0-9]+]],   [retval0+4];
+  ; CHECK:       ld.param.b32    [[R9:%r[0-9]+]],   [retval0+8];
+  ; CHECK:       ld.param.b32    [[R10:%r[0-9]+]],  [retval0+12];
+  ; CHECK:       ld.param.b32    [[R11:%r[0-9]+]],  [retval0+16];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R7]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R8]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R9]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R10]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R11]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St20 @callee_St20(%struct.St20* noundef byval(%struct.St20) align 4 %in)
+  %1 = extractvalue %struct.St20 %call, 0
+  %.fca.0.extract = extractvalue [5 x i32] %1, 0
+  %.fca.1.extract = extractvalue [5 x i32] %1, 1
+  %.fca.2.extract = extractvalue [5 x i32] %1, 2
+  %.fca.3.extract = extractvalue [5 x i32] %1, 3
+  %.fca.4.extract = extractvalue [5 x i32] %1, 4
+  %.fca.0.0.insert = insertvalue %struct.St20 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St20 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St20 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St20 %.fca.0.2.insert, i32 %.fca.3.extract, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St20 %.fca.0.3.insert, i32 %.fca.4.extract, 0, 4
+  ret %struct.St20 %.fca.0.4.insert
+}
+
+define internal %struct.St24 @callee_St24(%struct.St24* nocapture noundef readonly byval(%struct.St24) align 4 %in) {
+  ; CHECK:       .func  (.param .align 8 .b8 func_retval0[24])
+  ; CHECK-LABEL: callee_St24(
+  ; CHECK-NEXT:  .param .align 8 .b8 callee_St24_param_0[24]
+  ; CHECK:       ld.param.v2.u32 {[[R1:%r[0-9]+]],  [[R2:%r[0-9]+]]}, [callee_St24_param_0];
+  ; CHECK:       ld.param.v2.u32 {[[R3:%r[0-9]+]],  [[R4:%r[0-9]+]]}, [callee_St24_param_0+8];
+  ; CHECK:       ld.param.v2.u32 {[[R5:%r[0-9]+]],  [[R6:%r[0-9]+]]}, [callee_St24_param_0+16];
+  ; CHECK:       st.param.v2.b32 [func_retval0+0],  {[[R1]], [[R2]]};
+  ; CHECK:       st.param.v2.b32 [func_retval0+8],  {[[R3]], [[R4]]};
+  ; CHECK:       st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]};
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %arrayidx.4 = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 4
+  %5 = load i32, i32* %arrayidx.4, align 4
+  %arrayidx.5 = getelementptr inbounds %struct.St24, %struct.St24* %in, i32 0, i32 0, i32 5
+  %6 = load i32, i32* %arrayidx.5, align 4
+  %.fca.0.0.insert = insertvalue %struct.St24 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St24 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St24 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St24 %.fca.0.2.insert, i32 %4, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St24 %.fca.0.3.insert, i32 %5, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St24 %.fca.0.4.insert, i32 %6, 0, 5
+  ret %struct.St24 %.fca.0.5.insert
+}
+
+define %struct.St24 @caller_St24(%struct.St24* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[24])
+  ; CHECK-LABEL: caller_St24(
+  ; CHECK-NEXT:  .param .b32 caller_St24_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St24_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+12];
+  ; CHECK:       ld.u32          [[R5:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       ld.u32          [[R6:%r[0-9]+]], [[[R1]]+20];
+  ; CHECK:       ld.u32          [[R7:%r[0-9]+]], [[[R1]]+16];
+  ; CHECK:       .param .align 8 .b8 param0[24];
+  ; CHECK:       st.param.v2.b32 [param0+0],  {[[R3]], [[R2]]};
+  ; CHECK:       st.param.v2.b32 [param0+8],  {[[R5]], [[R4]]};
+  ; CHECK:       st.param.v2.b32 [param0+16], {[[R7]], [[R6]]};
+  ; CHECK:       .param .align 8 .b8 retval0[24];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St24, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.v2.b32 {[[R8:%r[0-9]+]],  [[R9:%r[0-9]+]]},  [retval0+0];
+  ; CHECK:       ld.param.v2.b32 {[[R10:%r[0-9]+]], [[R11:%r[0-9]+]]}, [retval0+8];
+  ; CHECK:       ld.param.v2.b32 {[[R12:%r[0-9]+]], [[R13:%r[0-9]+]]}, [retval0+16];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R8]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R9]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R10]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R11]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R12]];
+  ; CHECK:       st.param.b32    [func_retval0+20], [[R13]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St24 @callee_St24(%struct.St24* noundef byval(%struct.St24) align 4 %in)
+  %1 = extractvalue %struct.St24 %call, 0
+  %.fca.0.extract = extractvalue [6 x i32] %1, 0
+  %.fca.1.extract = extractvalue [6 x i32] %1, 1
+  %.fca.2.extract = extractvalue [6 x i32] %1, 2
+  %.fca.3.extract = extractvalue [6 x i32] %1, 3
+  %.fca.4.extract = extractvalue [6 x i32] %1, 4
+  %.fca.5.extract = extractvalue [6 x i32] %1, 5
+  %.fca.0.0.insert = insertvalue %struct.St24 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St24 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St24 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St24 %.fca.0.2.insert, i32 %.fca.3.extract, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St24 %.fca.0.3.insert, i32 %.fca.4.extract, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St24 %.fca.0.4.insert, i32 %.fca.5.extract, 0, 5
+  ret %struct.St24 %.fca.0.5.insert
+}
+
+define internal %struct.St28 @callee_St28(%struct.St28* nocapture noundef readonly byval(%struct.St28) align 4 %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[28])
+  ; CHECK-LABEL: callee_St28(
+  ; CHECK-NEXT:  .param .align 4 .b8 callee_St28_param_0[28]
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]],   [callee_St28_param_0];
+  ; CHECK:       ld.param.u32    [[R2:%r[0-9]+]],   [callee_St28_param_0+4];
+  ; CHECK:       ld.param.u32    [[R3:%r[0-9]+]],   [callee_St28_param_0+8];
+  ; CHECK:       ld.param.u32    [[R4:%r[0-9]+]],   [callee_St28_param_0+12];
+  ; CHECK:       ld.param.u32    [[R5:%r[0-9]+]],   [callee_St28_param_0+16];
+  ; CHECK:       ld.param.u32    [[R6:%r[0-9]+]],   [callee_St28_param_0+20];
+  ; CHECK:       ld.param.u32    [[R7:%r[0-9]+]],   [callee_St28_param_0+24];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R1]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R2]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R3]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R4]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R5]];
+  ; CHECK:       st.param.b32    [func_retval0+20], [[R6]];
+  ; CHECK:       st.param.b32    [func_retval0+24], [[R7]];
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %arrayidx.4 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 4
+  %5 = load i32, i32* %arrayidx.4, align 4
+  %arrayidx.5 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 5
+  %6 = load i32, i32* %arrayidx.5, align 4
+  %arrayidx.6 = getelementptr inbounds %struct.St28, %struct.St28* %in, i32 0, i32 0, i32 6
+  %7 = load i32, i32* %arrayidx.6, align 4
+  %.fca.0.0.insert = insertvalue %struct.St28 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St28 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St28 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St28 %.fca.0.2.insert, i32 %4, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St28 %.fca.0.3.insert, i32 %5, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St28 %.fca.0.4.insert, i32 %6, 0, 5
+  %.fca.0.6.insert = insertvalue %struct.St28 %.fca.0.5.insert, i32 %7, 0, 6
+  ret %struct.St28 %.fca.0.6.insert
+}
+
+define %struct.St28 @caller_St28(%struct.St28* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[28])
+  ; CHECK-LABEL: caller_St28(
+  ; CHECK-NEXT:  .param .b32 caller_St28_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St28_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       ld.u32          [[R5:%r[0-9]+]], [[[R1]]+12];
+  ; CHECK:       ld.u32          [[R6:%r[0-9]+]], [[[R1]]+16];
+  ; CHECK:       ld.u32          [[R7:%r[0-9]+]], [[[R1]]+20];
+  ; CHECK:       ld.u32          [[R8:%r[0-9]+]], [[[R1]]+24];
+  ; CHECK:       .param .align 4 .b8 param0[28];
+  ; CHECK:       st.param.b32    [param0+0],  [[R2]];
+  ; CHECK:       st.param.b32    [param0+4],  [[R3]];
+  ; CHECK:       st.param.b32    [param0+8],  [[R4]];
+  ; CHECK:       st.param.b32    [param0+12], [[R5]];
+  ; CHECK:       st.param.b32    [param0+16], [[R6]];
+  ; CHECK:       st.param.b32    [param0+20], [[R7]];
+  ; CHECK:       st.param.b32    [param0+24], [[R8]];
+  ; CHECK:       .param .align 4 .b8 retval0[28];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St28, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.b32    [[R9:%r[0-9]+]],   [retval0+0];
+  ; CHECK:       ld.param.b32    [[R10:%r[0-9]+]],  [retval0+4];
+  ; CHECK:       ld.param.b32    [[R11:%r[0-9]+]],  [retval0+8];
+  ; CHECK:       ld.param.b32    [[R12:%r[0-9]+]],  [retval0+12];
+  ; CHECK:       ld.param.b32    [[R13:%r[0-9]+]],  [retval0+16];
+  ; CHECK:       ld.param.b32    [[R14:%r[0-9]+]],  [retval0+20];
+  ; CHECK:       ld.param.b32    [[R15:%r[0-9]+]],  [retval0+24];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R9]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R10]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R11]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R12]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R13]];
+  ; CHECK:       st.param.b32    [func_retval0+20], [[R14]];
+  ; CHECK:       st.param.b32    [func_retval0+24], [[R15]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St28 @callee_St28(%struct.St28* noundef byval(%struct.St28) align 4 %in)
+  %1 = extractvalue %struct.St28 %call, 0
+  %.fca.0.extract = extractvalue [7 x i32] %1, 0
+  %.fca.1.extract = extractvalue [7 x i32] %1, 1
+  %.fca.2.extract = extractvalue [7 x i32] %1, 2
+  %.fca.3.extract = extractvalue [7 x i32] %1, 3
+  %.fca.4.extract = extractvalue [7 x i32] %1, 4
+  %.fca.5.extract = extractvalue [7 x i32] %1, 5
+  %.fca.6.extract = extractvalue [7 x i32] %1, 6
+  %.fca.0.0.insert = insertvalue %struct.St28 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St28 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St28 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St28 %.fca.0.2.insert, i32 %.fca.3.extract, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St28 %.fca.0.3.insert, i32 %.fca.4.extract, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St28 %.fca.0.4.insert, i32 %.fca.5.extract, 0, 5
+  %.fca.0.6.insert = insertvalue %struct.St28 %.fca.0.5.insert, i32 %.fca.6.extract, 0, 6
+  ret %struct.St28 %.fca.0.6.insert
+}
+
+define internal %struct.St32 @callee_St32(%struct.St32* nocapture noundef readonly byval(%struct.St32) align 4 %in) {
+  ; CHECK:       .func  (.param .align 16 .b8 func_retval0[32])
+  ; CHECK-LABEL: callee_St32(
+  ; CHECK-NEXT:  .param .align 16 .b8 callee_St32_param_0[32]
+  ; CHECK:       ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St32_param_0];
+  ; CHECK:       ld.param.v4.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]], [[R7:%r[0-9]+]], [[R8:%r[0-9]+]]}, [callee_St32_param_0+16];
+  ; CHECK:       st.param.v4.b32 [func_retval0+0],  {[[R1]], [[R2]], [[R3]], [[R4]]};
+  ; CHECK:       st.param.v4.b32 [func_retval0+16], {[[R5]], [[R6]], [[R7]], [[R8]]};
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %arrayidx.4 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 4
+  %5 = load i32, i32* %arrayidx.4, align 4
+  %arrayidx.5 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 5
+  %6 = load i32, i32* %arrayidx.5, align 4
+  %arrayidx.6 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 6
+  %7 = load i32, i32* %arrayidx.6, align 4
+  %arrayidx.7 = getelementptr inbounds %struct.St32, %struct.St32* %in, i32 0, i32 0, i32 7
+  %8 = load i32, i32* %arrayidx.7, align 4
+  %.fca.0.0.insert = insertvalue %struct.St32 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St32 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St32 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St32 %.fca.0.2.insert, i32 %4, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St32 %.fca.0.3.insert, i32 %5, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St32 %.fca.0.4.insert, i32 %6, 0, 5
+  %.fca.0.6.insert = insertvalue %struct.St32 %.fca.0.5.insert, i32 %7, 0, 6
+  %.fca.0.7.insert = insertvalue %struct.St32 %.fca.0.6.insert, i32 %8, 0, 7
+  ret %struct.St32 %.fca.0.7.insert
+}
+
+define %struct.St32 @caller_St32(%struct.St32* nocapture noundef readonly %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[32])
+  ; CHECK-LABEL: caller_St32(
+  ; CHECK-NEXT:  .param .b32 caller_St32_param_0
+  ; CHECK:       ld.param.u32    [[R1:%r[0-9]+]], [caller_St32_param_0];
+  ; CHECK:       ld.u32          [[R2:%r[0-9]+]], [[[R1]]+12];
+  ; CHECK:       ld.u32          [[R3:%r[0-9]+]], [[[R1]]+8];
+  ; CHECK:       ld.u32          [[R4:%r[0-9]+]], [[[R1]]+4];
+  ; CHECK:       ld.u32          [[R5:%r[0-9]+]], [[[R1]]];
+  ; CHECK:       ld.u32          [[R6:%r[0-9]+]], [[[R1]]+28];
+  ; CHECK:       ld.u32          [[R7:%r[0-9]+]], [[[R1]]+24];
+  ; CHECK:       ld.u32          [[R8:%r[0-9]+]], [[[R1]]+20];
+  ; CHECK:       ld.u32          [[R9:%r[0-9]+]], [[[R1]]+16];
+  ; CHECK:       .param .align 16 .b8 param0[32];
+  ; CHECK:       st.param.v4.b32 [param0+0],  {[[R5]], [[R4]], [[R3]], [[R2]]};
+  ; CHECK:       st.param.v4.b32 [param0+16], {[[R9]], [[R8]], [[R7]], [[R6]]};
+  ; CHECK:       .param .align 16 .b8 retval0[32];
+  ; CHECK:       call.uni (retval0),
+  ; CHECK-NEXT:  callee_St32, 
+  ; CHECK-NEXT:  (
+  ; CHECK-NEXT:  param0
+  ; CHECK-NEXT:  );
+  ; CHECK:       ld.param.v4.b32 {[[R10:%r[0-9]+]], [[R11:%r[0-9]+]], [[R12:%r[0-9]+]], [[R13:%r[0-9]+]]}, [retval0+0];
+  ; CHECK:       ld.param.v4.b32 {[[R14:%r[0-9]+]], [[R15:%r[0-9]+]], [[R16:%r[0-9]+]], [[R17:%r[0-9]+]]}, [retval0+16];
+  ; CHECK:       st.param.b32    [func_retval0+0],  [[R10]];
+  ; CHECK:       st.param.b32    [func_retval0+4],  [[R11]];
+  ; CHECK:       st.param.b32    [func_retval0+8],  [[R12]];
+  ; CHECK:       st.param.b32    [func_retval0+12], [[R13]];
+  ; CHECK:       st.param.b32    [func_retval0+16], [[R14]];
+  ; CHECK:       st.param.b32    [func_retval0+20], [[R15]];
+  ; CHECK:       st.param.b32    [func_retval0+24], [[R16]];
+  ; CHECK:       st.param.b32    [func_retval0+28], [[R17]];
+  ; CHECK-NEXT:  ret;
+  %call = tail call %struct.St32 @callee_St32(%struct.St32* noundef byval(%struct.St32) align 4 %in)
+  %1 = extractvalue %struct.St32 %call, 0
+  %.fca.0.extract = extractvalue [8 x i32] %1, 0
+  %.fca.1.extract = extractvalue [8 x i32] %1, 1
+  %.fca.2.extract = extractvalue [8 x i32] %1, 2
+  %.fca.3.extract = extractvalue [8 x i32] %1, 3
+  %.fca.4.extract = extractvalue [8 x i32] %1, 4
+  %.fca.5.extract = extractvalue [8 x i32] %1, 5
+  %.fca.6.extract = extractvalue [8 x i32] %1, 6
+  %.fca.7.extract = extractvalue [8 x i32] %1, 7
+  %.fca.0.0.insert = insertvalue %struct.St32 poison, i32 %.fca.0.extract, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St32 %.fca.0.0.insert, i32 %.fca.1.extract, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St32 %.fca.0.1.insert, i32 %.fca.2.extract, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St32 %.fca.0.2.insert, i32 %.fca.3.extract, 0, 3
+  %.fca.0.4.insert = insertvalue %struct.St32 %.fca.0.3.insert, i32 %.fca.4.extract, 0, 4
+  %.fca.0.5.insert = insertvalue %struct.St32 %.fca.0.4.insert, i32 %.fca.5.extract, 0, 5
+  %.fca.0.6.insert = insertvalue %struct.St32 %.fca.0.5.insert, i32 %.fca.6.extract, 0, 6
+  %.fca.0.7.insert = insertvalue %struct.St32 %.fca.0.6.insert, i32 %.fca.7.extract, 0, 7
+  ret %struct.St32 %.fca.0.7.insert
+}
+
+; Section 2 - checking that function argument (including retval) vectorization is done with private linkage;
+
+define private %struct.St16 @callee_St16_private(%struct.St16* nocapture noundef readonly byval(%struct.St16) align 4 %in) {
+  ; CHECK:       .func  (.param .align 16 .b8 func_retval0[16])
+  ; CHECK-LABEL: callee_St16_private(
+  ; CHECK-NEXT:  .param .align 16 .b8 callee_St16_private_param_0[16]
+  ; CHECK:       ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St16_private_param_0];
+  ; CHECK:       st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %.fca.0.0.insert = insertvalue %struct.St16 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St16 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St16 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St16 %.fca.0.2.insert, i32 %4, 0, 3
+  ret %struct.St16 %.fca.0.3.insert
+}
+
+; Section 3 - checking that function argument (including retval) vectorization 
+; is NOT done with linkage types other than internal and private.
+
+define external %struct.St16 @callee_St16_external(%struct.St16* nocapture noundef readonly byval(%struct.St16) align 4 %in) {
+  ; CHECK:       .func  (.param .align 4 .b8 func_retval0[16])
+  ; CHECK-LABEL: callee_St16_external(
+  ; CHECK-NEXT:  .param .align 4 .b8 callee_St16_external_param_0[16]
+  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]],   [callee_St16_external_param_0];
+  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]],   [callee_St16_external_param_0+4];
+  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]],   [callee_St16_external_param_0+8];
+  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]],   [callee_St16_external_param_0+12];
+  ; CHECK:       st.param.b32 [func_retval0+0],  [[R1]];
+  ; CHECK:       st.param.b32 [func_retval0+4],  [[R2]];
+  ; CHECK:       st.param.b32 [func_retval0+8],  [[R3]];
+  ; CHECK:       st.param.b32 [func_retval0+12], [[R4]];
+  ; CHECK-NEXT:  ret;
+  %arrayidx = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 0
+  %1 = load i32, i32* %arrayidx, align 4
+  %arrayidx.1 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 1
+  %2 = load i32, i32* %arrayidx.1, align 4
+  %arrayidx.2 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 2
+  %3 = load i32, i32* %arrayidx.2, align 4
+  %arrayidx.3 = getelementptr inbounds %struct.St16, %struct.St16* %in, i32 0, i32 0, i32 3
+  %4 = load i32, i32* %arrayidx.3, align 4
+  %.fca.0.0.insert = insertvalue %struct.St16 poison, i32 %1, 0, 0
+  %.fca.0.1.insert = insertvalue %struct.St16 %.fca.0.0.insert, i32 %2, 0, 1
+  %.fca.0.2.insert = insertvalue %struct.St16 %.fca.0.1.insert, i32 %3, 0, 2
+  %.fca.0.3.insert = insertvalue %struct.St16 %.fca.0.2.insert, i32 %4, 0, 3
+  ret %struct.St16 %.fca.0.3.insert
+}