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-vectors.ll

@@ -353,6 +353,76 @@ define void @global_2xf32_no_align(ptr addrspace(1) %a, ptr addrspace(1) %b) {
   ret void
 }
 
+; 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<13>;
+; CHECK-NEXT:    .reg .b32 %r<20>;
+; 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:    ld.b16 %rs7, [%rd1+16];
+; CHECK-NEXT:    mov.b32 %r12, {%rs7, %rs8};
+; CHECK-NEXT:    ld.b32 %r13, [%rd1+12];
+; CHECK-NEXT:    max.f16x2 %r14, %r12, %r7;
+; CHECK-NEXT:    max.f16x2 %r15, %r13, %r7;
+; CHECK-NEXT:    st.b32 [%rd1+12], %r15;
+; CHECK-NEXT:    mov.b32 {%rs9, _}, %r14;
+; CHECK-NEXT:    st.b16 [%rd1+16], %rs9;
+; CHECK-NEXT:    ld.b16 %rs10, [%rd1+22];
+; CHECK-NEXT:    mov.b32 %r16, {%rs10, %rs11};
+; CHECK-NEXT:    ld.b32 %r17, [%rd1+18];
+; CHECK-NEXT:    max.f16x2 %r18, %r16, %r7;
+; CHECK-NEXT:    max.f16x2 %r19, %r17, %r7;
+; CHECK-NEXT:    st.b32 [%rd1+18], %r19;
+; CHECK-NEXT:    mov.b32 {%rs12, _}, %r18;
+; CHECK-NEXT:    st.b16 [%rd1+22], %rs12;
+; CHECK-NEXT:    ret;
+  %1 = call <8 x half> @llvm.masked.load.v8f16.p0(ptr align 16 %rd0, <8 x i1> <i1 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 false, i1 false>, <8 x half> poison)
+  %load13 = shufflevector <8 x half> %1, <8 x half> poison, <3 x i32> <i32 0, i32 1, i32 2>
+  %load24 = shufflevector <8 x half> %1, <8 x half> poison, <3 x i32> <i32 3, i32 4, i32 5>
+  %Extend5 = extractelement <8 x half> %1, i32 6
+  %Extend26 = extractelement <8 x half> %1, i32 7
+  %p1 = fcmp ogt <3 x half> %load13, zeroinitializer
+  %s1 = select <3 x i1> %p1, <3 x half> %load13, <3 x half> zeroinitializer
+  store <3 x half> %s1, ptr %rd0, align 16
+  %in2 = getelementptr half, ptr %rd0, i64 3
+  %p2 = fcmp ogt <3 x half> %load24, zeroinitializer
+  %s2 = select <3 x i1> %p2, <3 x half> %load24, <3 x half> zeroinitializer
+  store <3 x half> %s2, ptr %in2, align 4
+  %in3 = getelementptr half, ptr %rd0, i64 6
+  %load3 = load <3 x half>, ptr %in3, align 4
+  %p3 = fcmp ogt <3 x half> %load3, zeroinitializer
+  %s3 = select <3 x i1> %p3, <3 x half> %load3, <3 x half> zeroinitializer
+  store <3 x half> %s3, ptr %in3, align 4
+  %in4 = getelementptr half, ptr %rd0, i64 9
+  %load4 = load <3 x half>, ptr %in4, align 4
+  %p4 = fcmp ogt <3 x half> %load4, zeroinitializer
+  %s4 = select <3 x i1> %p4, <3 x half> %load4, <3 x half> zeroinitializer
+  store <3 x half> %s4, ptr %in4, align 4
+  ret void
+}
+
 declare <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1), <8 x i1>, <8 x i32>)
 declare void @llvm.masked.store.v8i32.p1(<8 x i32>, ptr addrspace(1), <8 x i1>)
 declare <16 x i16> @llvm.masked.load.v16i16.p1(ptr addrspace(1), <16 x i1>, <16 x i16>)
@@ -363,4 +433,5 @@ declare <4 x i8> @llvm.masked.load.v4i8.p1(ptr addrspace(1), <4 x i1>, <4 x i8>)
 declare void @llvm.masked.store.v4i8.p1(<4 x i8>, ptr addrspace(1), <4 x i1>)
 declare <2 x float> @llvm.masked.load.v2f32.p1(ptr addrspace(1), <2 x i1>, <2 x float>)
 declare void @llvm.masked.store.v2f32.p1(<2 x float>, ptr addrspace(1), <2 x i1>)
+declare <8 x half> @llvm.masked.load.v8f16.p0(ptr captures(none), <8 x i1>, <8 x half>)
 !0 = !{}