Add new test, adjust comments

Author: dakersnar

llvm/lib/Transforms/Vectorize/LoadStoreVectorizer.cpp

@@ -351,8 +351,10 @@ class Vectorizer {
   /// in the chain is the leader, and an instr touches distance 0 from itself.
   std::vector<Chain> gatherChains(ArrayRef<Instruction *> Instrs);
 
-  /// Is a load/store with this alignment allowed by TTI and at least as fast
-  /// as an unvectorized load/store.
+  /// Checks if a potential vector load/store with a given alignment is allowed
+  /// and fast. Aligned accesses are always allowed and fast, while misaligned
+  /// accesses depend on TTI checks to determine whether they can and should be
+  /// vectorized or kept as element-wise accesses.
   bool accessIsAllowedAndFast(unsigned SizeBytes, unsigned AS, Align Alignment,
                               unsigned VecElemBits) const;
 
@@ -1909,7 +1911,7 @@ bool Vectorizer::accessIsAllowedAndFast(unsigned SizeBytes, unsigned AS,
   if (Alignment.value() % SizeBytes == 0)
     return true;
 
-  // Element-wise access *might* be faster than misaligned vector accesses.
+  // Ask TTI whether misaligned accesses are faster as vector or element-wise.
   unsigned VectorizedSpeed = 0;
   bool AllowsMisaligned = TTI.allowsMisalignedMemoryAccesses(
       F.getContext(), SizeBytes * 8, AS, Alignment, &VectorizedSpeed);

llvm/test/CodeGen/NVPTX/masked-load-3xhalf.ll

@@ -0,0 +1,84 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}
+
+; This is testing the lowering behavior of this case from LoadStoreVectorizer/NVPTX/4x2xhalf.ll
+; where two 3xhalfs are chained together and extended to 8xhalf.
+define void @halfx3_extend_chain(ptr align 16 captures(none) %rd0) {
+; CHECK-LABEL: halfx3_extend_chain(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<7>;
+; CHECK-NEXT:    .reg .b32 %r<12>;
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [halfx3_extend_chain_param_0];
+; CHECK-NEXT:    .pragma "used_bytes_mask 0xfff";
+; CHECK-NEXT:    ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; CHECK-NEXT:    mov.b32 {%rs1, %rs2}, %r3;
+; CHECK-NEXT:    mov.b32 {_, %rs3}, %r2;
+; CHECK-NEXT:    mov.b32 %r5, {%rs3, %rs1};
+; CHECK-NEXT:    mov.b32 %r6, {%rs2, %rs4};
+; CHECK-NEXT:    mov.b32 %r7, 0;
+; CHECK-NEXT:    max.f16x2 %r8, %r2, %r7;
+; CHECK-NEXT:    max.f16x2 %r9, %r1, %r7;
+; CHECK-NEXT:    st.b32 [%rd1], %r9;
+; CHECK-NEXT:    mov.b32 {%rs5, _}, %r8;
+; CHECK-NEXT:    st.b16 [%rd1+4], %rs5;
+; CHECK-NEXT:    max.f16x2 %r10, %r6, %r7;
+; CHECK-NEXT:    max.f16x2 %r11, %r5, %r7;
+; CHECK-NEXT:    st.b32 [%rd1+6], %r11;
+; CHECK-NEXT:    mov.b32 {%rs6, _}, %r10;
+; CHECK-NEXT:    st.b16 [%rd1+10], %rs6;
+; CHECK-NEXT:    ret;
+  %load1 = load <3 x half>, ptr %rd0, align 16
+  %p1 = fcmp ogt <3 x half> %load1, zeroinitializer
+  %s1 = select <3 x i1> %p1, <3 x half> %load1, <3 x half> zeroinitializer
+  store <3 x half> %s1, ptr %rd0, align 16
+  %in2 = getelementptr half, ptr %rd0, i64 3
+  %load2 = load <3 x half>, ptr %in2, align 4
+  %p2 = fcmp ogt <3 x half> %load2, zeroinitializer
+  %s2 = select <3 x i1> %p2, <3 x half> %load2, <3 x half> zeroinitializer
+  store <3 x half> %s2, ptr %in2, align 4
+  ret void
+}
+
+; This disables the vectorization by reducing the alignment.
+define void @halfx3_no_align(ptr align 4 captures(none) %rd0) {
+; CHECK-LABEL: halfx3_no_align(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<7>;
+; CHECK-NEXT:    .reg .b32 %r<10>;
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [halfx3_no_align_param_0];
+; CHECK-NEXT:    ld.b16 %rs1, [%rd1+4];
+; CHECK-NEXT:    mov.b32 %r1, {%rs1, %rs2};
+; CHECK-NEXT:    ld.b32 %r2, [%rd1];
+; CHECK-NEXT:    mov.b32 %r3, 0;
+; CHECK-NEXT:    max.f16x2 %r4, %r1, %r3;
+; CHECK-NEXT:    max.f16x2 %r5, %r2, %r3;
+; CHECK-NEXT:    st.b32 [%rd1], %r5;
+; CHECK-NEXT:    mov.b32 {%rs3, _}, %r4;
+; CHECK-NEXT:    st.b16 [%rd1+4], %rs3;
+; CHECK-NEXT:    ld.b16 %rs4, [%rd1+10];
+; CHECK-NEXT:    mov.b32 %r6, {%rs4, %rs5};
+; CHECK-NEXT:    ld.b32 %r7, [%rd1+6];
+; CHECK-NEXT:    max.f16x2 %r8, %r6, %r3;
+; CHECK-NEXT:    max.f16x2 %r9, %r7, %r3;
+; CHECK-NEXT:    st.b32 [%rd1+6], %r9;
+; CHECK-NEXT:    mov.b32 {%rs6, _}, %r8;
+; CHECK-NEXT:    st.b16 [%rd1+10], %rs6;
+; CHECK-NEXT:    ret;
+  %load1 = load <3 x half>, ptr %rd0, align 4
+  %p1 = fcmp ogt <3 x half> %load1, zeroinitializer
+  %s1 = select <3 x i1> %p1, <3 x half> %load1, <3 x half> zeroinitializer
+  store <3 x half> %s1, ptr %rd0, align 4
+  %in2 = getelementptr half, ptr %rd0, i64 3
+  %load2 = load <3 x half>, ptr %in2, align 4
+  %p2 = fcmp ogt <3 x half> %load2, zeroinitializer
+  %s2 = select <3 x i1> %p2, <3 x half> %load2, <3 x half> zeroinitializer
+  store <3 x half> %s2, ptr %in2, align 4
+  ret void
+}