[NVPTX] Unforce minimum alignment of 4 for byval arguments of device-side functions.

[pavelkopyl]

Summary

Minimum alignment of 4 for byval arguments was forced to workaround
a bug in old versions of ptxas. Details: https://reviews.llvm.org/D22428.
Recent ptxas versions (> 9.0) do not seem to have this bug, so alignment
requirement was relaxed. To force again minimum alignment of 4, use
'-force-min-byval-param-align' option.

[tra]

Do we need to do it? I mean -- are there any observable differences between enforced alignment of 4 and not? For register-passed parameters it will make no difference, and for the parameters passed via constant/local memory, it may or may not make a difference.

I guess at the extreme it would impact how many 1-byte parameters we may have, as the total size of param address space has an upper bound (for kernels it's 4K: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#parameter-buffer-layout). If that's the case, then it probably does not matter whether we can pass 4K arguments or only 1K.

In other words, I do not see a practical need to introduce an option which may not have any practical use.

[pavelkopyl]

In D141737#4053125, @tra wrote:

Do we need to do it? I mean -- are there any observable differences between enforced alignment of 4 and not? For register-passed parameters it will make no difference, and for the parameters passed via constant/local memory, it may or may not make a difference.

I guess at the extreme it would impact how many 1-byte parameters we may have, as the total size of param address space has an upper bound (for kernels it's 4K: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#parameter-buffer-layout). If that's the case, then it probably does not matter whether we can pass 4K arguments or only 1K.

In other words, I do not see a practical need to introduce an option which may not have any practical use.

Sorry for long response.
OK, I'll remove the option. I just wanted to clarify, are you agree in general with removing this workaround?

[pavelkopyl]

I guess, there is one more reason to getting back to default alignment.It seems we break ABI with alignment forced to 4. This may be an issue for mixed build clang + nvcc.

[pavelkopyl]

Hi @tra, does it make sense ABI related reason for you, or it's reasonable just to drop this change?

[pavelkopyl]

-Rebase

[tra]

In D141737#4292014, @pavelkopyl wrote:

Hi @tra, does it make sense ABI related reason for you, or it's reasonable just to drop this change?

Yes. Removing this workaround makes sense.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

cl::desc("NVPTX Specifies: 0 use div.approx, 1 use div.full, 2 use"
              " IEEE Compliant F32 div.rnd if available."),
     cl::init(2));
 
 static cl::opt<bool> UsePrecSqrtF32(
     "nvptx-prec-sqrtf32", cl::Hidden,
     cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."),
     cl::init(true));
 
+static cl::opt<bool> ForceMinByValParamAlign(
+    "nvptx-force-min-byval-param-align", cl::Hidden,
+    cl::desc("NVPTX Specific: force 4-byte minimal alignment for byval"
+             " params of device functions."),
+    cl::init(false));
+
 int NVPTXTargetLowering::getDivF32Level() const {
   if (UsePrecDivF32.getNumOccurrences() > 0) {
     // If nvptx-prec-div32=N is used on the command-line, always honor it
     return UsePrecDivF32;
   } else {
     // Otherwise, use div.approx if fast math is enabled
     if (getTargetMachine().Options.UnsafeFPMath)
       return 0;

 Align NVPTXTargetLowering::getFunctionByValParamAlign(
     const Function *F, Type *ArgTy, Align InitialAlign,
     const DataLayout &DL) const {
   Align ArgAlign = InitialAlign;
   // Try to increase alignment to enhance vectorization options.
   if (F)
     ArgAlign = std::max(ArgAlign, getFunctionParamOptimizedAlign(F, ArgTy, DL));
 
-  // Work around a bug in ptxas. When PTX code takes address of
+  // Old ptx versions have a bug. 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.
-  // 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.
+  // least 4. This bug seems to be fixed at least starting from
+  // ptxas > 9.0.
+  // TODO: remove this after verifying the bug is not reproduced
+  // on non-deprecated ptxas versions.
+  if (ForceMinByValParamAlign)
     ArgAlign = std::max(ArgAlign, Align(4));
 
   return ArgAlign;
 }
 
 /// 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

llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll

 ; RUN: llc < %s -march=nvptx -mcpu=sm_50 -verify-machineinstrs | FileCheck %s
 ; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_50 -verify-machineinstrs | %ptxas-verify %}
 
 ; calls with a bitcasted function symbol should be fine, but in combination with
 ; a byval attribute were causing a segfault during isel. This testcase was
 ; reduced from a SYCL kernel using aggregate types which ended up being passed
 ; `byval`
 
 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 %"class.complex" = type { %"class.sycl::_V1::detail::half_impl::half", %"class.sycl::_V1::detail::half_impl::half" }
 %"class.sycl::_V1::detail::half_impl::half" = type { half }
 %complex_half = type { half, half }
 
-; CHECK: .param .align 4 .b8 param2[4];
-; CHECK: st.param.v2.b16         [param2+0], {%h2, %h1};
+; CHECK: .param .align 2 .b8 param2[4];
+; CHECK: st.param.b16   [param2+0], %h1;
+; CHECK: st.param.b16   [param2+2], %h2;
 ; CHECK: .param .align 2 .b8 retval0[4];
 ; CHECK: call.uni (retval0),
 ; CHECK-NEXT: _Z20__spirv_GroupCMulKHRjjN5__spv12complex_halfE,
 define weak_odr void @foo() {
 entry:
   %call.i.i.i = tail call %"class.complex" @_Z20__spirv_GroupCMulKHRjjN5__spv12complex_halfE(i32 0, i32 0, ptr byval(%"class.complex") null)
   ret void
 }
 
 ;; Function pointers can escape, so we have to use a conservative
 ;; alignment for a function that has address taken.
 ;;
 declare ptr @usefp(ptr %fp)
 ; CHECK: .func callee(
-; CHECK-NEXT: .param .align 4 .b8 callee_param_0[4]
+; CHECK-NEXT: .param .align 2 .b8 callee_param_0[4]
 define internal void @callee(ptr byval(%"class.complex") %byval_arg) {
   ret void
 }
 define void @boom() {
   %fp = call ptr @usefp(ptr @callee)
-  ; CHECK: .param .align 4 .b8 param0[4];
-  ; CHECK: st.param.v2.b16 [param0+0]
-  ; CHECK: .callprototype ()_ (.param .align 4 .b8 _[4]);
+  ; CHECK: .param .align 2 .b8 param0[4];
+  ; CHECK: st.param.b16 [param0+0], %h1;
+  ; CHECK: st.param.b16 [param0+2], %h2;
+  ; CHECK: .callprototype ()_ (.param .align 2 .b8 _[4]);
   call void %fp(ptr byval(%"class.complex") null)
   ret void
 }
 
 declare %complex_half @_Z20__spirv_GroupCMulKHRjjN5__spv12complex_halfE(i32, i32, ptr byval(%"class.complex"))

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

-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=CHECK,NOALIGN4
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-force-min-byval-param-align | FileCheck %s --check-prefixes=CHECK,ALIGN4
 ; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-force-min-byval-param-align | %ptxas-verify %}
 
 ;;; Need 4-byte alignment on ptr passed byval
 define ptx_device void @t1(ptr byval(float) %x) {
 ; CHECK: .func t1
 ; CHECK: .param .align 4 .b8 t1_param_0[4]
   ret void
 }
 

 ;;; Need 4-byte alignment on float2* passed byval
 %struct.float2 = type { float, float }
 define ptx_device void @t3(ptr byval(%struct.float2) %x) {
 ; CHECK: .func t3
 ; CHECK: .param .align 4 .b8 t3_param_0[8]
   ret void
 }
 
-;;; Need at least 4-byte alignment in order to avoid miscompilation by
-;;; ptxas for sm_50+
 define ptx_device void @t4(ptr byval(i8) %x) {
 ; CHECK: .func t4
-; CHECK: .param .align 4 .b8 t4_param_0[1]
+; NOALIGN4: .param .align 1 .b8 t4_param_0[1]
+; ALIGN4: .param .align 4 .b8 t4_param_0[1]
   ret void
 }
 
 ;;; Make sure we adjust alignment at the call site as well.
 define ptx_device void @t5(ptr align 2 byval(i8) %x) {
 ; CHECK: .func t5
-; CHECK: .param .align 4 .b8 t5_param_0[1]
+; NOALIGN4: .param .align 2 .b8 t5_param_0[1]
+; ALIGN4: .param .align 4 .b8 t5_param_0[1]
 ; CHECK: {
-; CHECK: .param .align 4 .b8 param0[1];
+; NOALIGN4: .param .align 1 .b8 param0[1];
+; ALIGN4:   .param .align 4 .b8 param0[1];
 ; CHECK: call.uni
   call void @t4(ptr byval(i8) %x)
   ret void
 }
 
 ;;; Make sure we adjust alignment for a function prototype
 ;;; in case of an inderect call.
 
 declare ptr @getfp(i32 %n)
 %struct.half2 = type { half, half }
 define ptx_device void @t6() {
 ; CHECK: .func t6
   %fp = call ptr @getfp(i32 0)
 ; CHECK: prototype_2 : .callprototype ()_ (.param .align 8 .b8 _[8]);
   call void %fp(ptr byval(double) null);
 
   %fp2 = call ptr @getfp(i32 1)
-; CHECK: prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
+; NOALIGN4: prototype_4 : .callprototype ()_ (.param .align 2 .b8 _[4]);
+; ALIGN4: prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
   call void %fp(ptr byval(%struct.half2) null);
 
   %fp3 = call ptr @getfp(i32 2)
-; CHECK: prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
+; NOALIGN4: prototype_6 : .callprototype ()_ (.param .align 1 .b8 _[1]);
+; ALIGN4: prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
   call void %fp(ptr byval(i8) null);
   ret void
 }