Fixed sm version for .and bmma operator.

[JackAKirk]

Summary

As stated here (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-mma):

".and operation in single-bit wmma requires sm_80 or higher."

[tra]

Good catch. LGTM.

[JackAKirk]

In D131265#3702594, @tra wrote:

Good catch. LGTM.

Thanks. If you could land it for me that would be much appreciated. I don't have the rights.

[tra]

Looks like the tests needed to be updated (and I've found one bug which explains how we've missed this).

[JackAKirk]

In D131265#3702954, @tra wrote:

Looks like the tests needed to be updated (and I've found one bug which explains how we've missed this).

Ah yes I see it. Thanks for updating the tests.

clang/include/clang/Basic/BuiltinsNVPTX.def

 TARGET_BUILTIN(__hmma_m8n32k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", AND(SM_70,PTX61))
 TARGET_BUILTIN(__hmma_m8n32k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", AND(SM_70,PTX61))
 TARGET_BUILTIN(__hmma_m8n32k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", AND(SM_70,PTX61))
 
 // Builtins to support integer and sub-integer WMMA instructions on sm_72/sm_75
 TARGET_BUILTIN(__bmma_m8n8k128_ld_a_b1, "vi*iC*UiIi", "", AND(SM_75,PTX63))
 TARGET_BUILTIN(__bmma_m8n8k128_ld_b_b1, "vi*iC*UiIi", "", AND(SM_75,PTX63))
 TARGET_BUILTIN(__bmma_m8n8k128_ld_c, "vi*iC*UiIi", "", AND(SM_75,PTX63))
-TARGET_BUILTIN(__bmma_m8n8k128_mma_and_popc_b1, "vi*iC*iC*iC*Ii", "", AND(SM_75,PTX71))
+TARGET_BUILTIN(__bmma_m8n8k128_mma_and_popc_b1, "vi*iC*iC*iC*Ii", "", AND(SM_80,PTX71))
 TARGET_BUILTIN(__bmma_m8n8k128_mma_xor_popc_b1, "vi*iC*iC*iC*Ii", "", AND(SM_75,PTX63))
 TARGET_BUILTIN(__bmma_m8n8k128_st_c_i32, "vi*iC*UiIi", "", AND(SM_75,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_ld_a_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_ld_a_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_ld_b_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_ld_b_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_ld_c, "vi*iC*UiIi", "", AND(SM_72,PTX63))
 TARGET_BUILTIN(__imma_m16n16k16_mma_s8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63))

clang/test/CodeGen/builtins-nvptx-mma.cu

 
 //
 // *** DO NOT EDIT ***
 //
 //  This test has been automatically generated by
 //  builtins-nvtx-mma.py --ptx=71 --gpu-arch=80
 //
 // Make sure we can handle all builtins available on sm_80 with PTX71
 // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \
 // RUN:            -fcuda-is-device -target-feature +ptx71 \
 // RUN:            -DPTX=71 -DSM=80 \
 // RUN:            -S -emit-llvm -o - -x cuda %s \
-// RUN:   | FileCheck -check-prefixes=CHECK_PTX70_SM80,CHECK_PTX60_SM70,CHECK_PTX63_SM72,CHECK_PTX61_SM70,CHECK_PTX63_SM75,CHECK_PTX71_SM75 %s
+// RUN:   | FileCheck -check-prefixes=CHECK_PTX70_SM80,CHECK_PTX60_SM70,CHECK_PTX63_SM72,CHECK_PTX61_SM70,CHECK_PTX63_SM75,CHECK_PTX71_SM80 %s
 // Verify that all builtins have correct constraints.
 // RUN: %clang_cc1 -triple nvptx-unknown-unknown \
 // RUN:   -target-cpu sm_60 -target-feature +ptx42 \
 // RUN:   -DPTX=71 -DSM=80 -fcuda-is-device -S -o /dev/null -x cuda \
 // RUN:   -verify %s
 
 
 #if !defined(CUDA_VERSION)
#if (PTX >= 60) && (SM >= 70)
   // expected-error-re@+1 {{'__hmma_m16n16k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx60{{.*}}}}
   __hmma_m16n16k16_mma_f32f32(fdst, src, src, fsrc, 1, 1);
   // CHECK_PTX60_SM70: call {{.*}} @llvm.nvvm.wmma.m16n16k16.mma.row.row.f32.f32
   // expected-error-re@+1 {{'__hmma_m16n16k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx60{{.*}}}}
   __hmma_m16n16k16_mma_f32f32(fdst, src, src, fsrc, 0, 0);
   // CHECK_PTX60_SM70: call {{.*}} @llvm.nvvm.wmma.m16n16k16.mma.row.row.f32.f32.satfinite
   // expected-error-re@+1 {{'__hmma_m16n16k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx60{{.*}}}}
   __hmma_m16n16k16_mma_f32f32(fdst, src, src, fsrc, 0, 1);
-#endif // (PTX >= 60) && (SM >= 70) 
+#endif // (PTX >= 60) && (SM >= 70)
 
 #if (PTX >= 61) && (SM >= 70)
 
   // CHECK_PTX61_SM70: call {{.*}} @llvm.nvvm.wmma.m32n8k16.load.a.col.stride.f16
   // expected-error-re@+1 {{'__hmma_m32n8k16_ld_a' needs target feature (sm_70{{.*}},(ptx61{{.*}}}}
   __hmma_m32n8k16_ld_a(dst, src, ldm, 1);
   // CHECK_PTX61_SM70: call {{.*}} @llvm.nvvm.wmma.m32n8k16.load.a.row.stride.f16
   // expected-error-re@+1 {{'__hmma_m32n8k16_ld_a' needs target feature (sm_70{{.*}},(ptx61{{.*}}}}
#if (PTX >= 61) && (SM >= 70)
   // expected-error-re@+1 {{'__hmma_m8n32k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx61{{.*}}}}
   __hmma_m8n32k16_mma_f32f32(fdst, src, src, fsrc, 1, 1);
   // CHECK_PTX61_SM70: call {{.*}} @llvm.nvvm.wmma.m8n32k16.mma.row.row.f32.f32
   // expected-error-re@+1 {{'__hmma_m8n32k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx61{{.*}}}}
   __hmma_m8n32k16_mma_f32f32(fdst, src, src, fsrc, 0, 0);
   // CHECK_PTX61_SM70: call {{.*}} @llvm.nvvm.wmma.m8n32k16.mma.row.row.f32.f32.satfinite
   // expected-error-re@+1 {{'__hmma_m8n32k16_mma_f32f32' needs target feature (sm_70{{.*}},(ptx61{{.*}}}}
   __hmma_m8n32k16_mma_f32f32(fdst, src, src, fsrc, 0, 1);
-#endif // (PTX >= 61) && (SM >= 70) 
+#endif // (PTX >= 61) && (SM >= 70)
 
 #if (PTX >= 63) && (SM >= 72)
 
   // CHECK_PTX63_SM72: call {{.*}} @llvm.nvvm.wmma.m16n16k16.load.a.col.stride.s8
   // expected-error-re@+1 {{'__imma_m16n16k16_ld_a_s8' needs target feature (sm_72{{.*}},(ptx63{{.*}}}}
   __imma_m16n16k16_ld_a_s8(dst, src, ldm, 1);
   // CHECK_PTX63_SM72: call {{.*}} @llvm.nvvm.wmma.m16n16k16.load.a.row.stride.s8
   // expected-error-re@+1 {{'__imma_m16n16k16_ld_a_s8' needs target feature (sm_72{{.*}},(ptx63{{.*}}}}
#if (PTX >= 63) && (SM >= 72)
   // expected-error-re@+1 {{'__imma_m8n32k16_mma_u8' needs target feature (sm_72{{.*}},(ptx63{{.*}}}}
   __imma_m8n32k16_mma_u8(dst, src, src, src, 1, 1);
   // CHECK_PTX63_SM72: call {{.*}} @llvm.nvvm.wmma.m8n32k16.mma.row.row.u8
   // expected-error-re@+1 {{'__imma_m8n32k16_mma_u8' needs target feature (sm_72{{.*}},(ptx63{{.*}}}}
   __imma_m8n32k16_mma_u8(dst, src, src, src, 0, 0);
   // CHECK_PTX63_SM72: call {{.*}} @llvm.nvvm.wmma.m8n32k16.mma.row.row.u8.satfinite
   // expected-error-re@+1 {{'__imma_m8n32k16_mma_u8' needs target feature (sm_72{{.*}},(ptx63{{.*}}}}
   __imma_m8n32k16_mma_u8(dst, src, src, src, 0, 1);
-#endif // (PTX >= 63) && (SM >= 72) 
+#endif // (PTX >= 63) && (SM >= 72)
 
 #if (PTX >= 63) && (SM >= 75)
 
   // CHECK_PTX63_SM75: call {{.*}} @llvm.nvvm.wmma.m8n8k128.load.a.row.stride.b1
   // expected-error-re@+1 {{'__bmma_m8n8k128_ld_a_b1' needs target feature (sm_75{{.*}},(ptx63{{.*}}}}
   __bmma_m8n8k128_ld_a_b1(dst, src, ldm, 0);
   // CHECK_PTX63_SM75: call {{.*}} @llvm.nvvm.wmma.m8n8k128.load.b.col.stride.b1
   // expected-error-re@+1 {{'__bmma_m8n8k128_ld_b_b1' needs target feature (sm_75{{.*}},(ptx63{{.*}}}}
#if (PTX >= 63) && (SM >= 75)
   // expected-error-re@+1 {{'__imma_m8n8k32_mma_s4' needs target feature (sm_75{{.*}},(ptx63{{.*}}}}
   __imma_m8n8k32_mma_s4(dst, src, src, src, 1, 1);
   // CHECK_PTX63_SM75: call {{.*}} @llvm.nvvm.wmma.m8n8k32.mma.row.col.u4
   // expected-error-re@+1 {{'__imma_m8n8k32_mma_u4' needs target feature (sm_75{{.*}},(ptx63{{.*}}}}
   __imma_m8n8k32_mma_u4(dst, src, src, src, 1, 0);
   // CHECK_PTX63_SM75: call {{.*}} @llvm.nvvm.wmma.m8n8k32.mma.row.col.u4.satfinite
   // expected-error-re@+1 {{'__imma_m8n8k32_mma_u4' needs target feature (sm_75{{.*}},(ptx63{{.*}}}}
   __imma_m8n8k32_mma_u4(dst, src, src, src, 1, 1);
-#endif // (PTX >= 63) && (SM >= 75) 
+#endif // (PTX >= 63) && (SM >= 75)
 
 #if (PTX >= 70) && (SM >= 80)
 
   // CHECK_PTX70_SM80: call {{.*}} @llvm.nvvm.wmma.m16n16k16.load.a.col.stride.bf16
   // expected-error-re@+1 {{'__mma_bf16_m16n16k16_ld_a' needs target feature (sm_80{{.*}},(ptx70{{.*}}}}
   __mma_bf16_m16n16k16_ld_a(dst, src, ldm, 1);
   // CHECK_PTX70_SM80: call {{.*}} @llvm.nvvm.wmma.m16n16k16.load.a.row.stride.bf16
   // expected-error-re@+1 {{'__mma_bf16_m16n16k16_ld_a' needs target feature (sm_80{{.*}},(ptx70{{.*}}}}
#if (PTX >= 70) && (SM >= 80)
   // expected-error-re@+1 {{'__dmma_m8n8k4_mma_f64' needs target feature (sm_80{{.*}},(ptx70{{.*}}}}
   __dmma_m8n8k4_mma_f64(ddst, dsrc, dsrc, dsrc, 2, 0);
   // CHECK_PTX70_SM80: call {{.*}} @llvm.nvvm.wmma.m8n8k4.mma.row.col.f64
   // expected-error-re@+1 {{'__dmma_m8n8k4_mma_f64' needs target feature (sm_80{{.*}},(ptx70{{.*}}}}
   __dmma_m8n8k4_mma_f64(ddst, dsrc, dsrc, dsrc, 1, 0);
   // CHECK_PTX70_SM80: call {{.*}} @llvm.nvvm.wmma.m8n8k4.mma.row.row.f64
   // expected-error-re@+1 {{'__dmma_m8n8k4_mma_f64' needs target feature (sm_80{{.*}},(ptx70{{.*}}}}
   __dmma_m8n8k4_mma_f64(ddst, dsrc, dsrc, dsrc, 0, 0);
-#endif // (PTX >= 70) && (SM >= 80) 
+#endif // (PTX >= 70) && (SM >= 80)
 
-#if (PTX >= 71) && (SM >= 75)
+#if (PTX >= 71) && (SM >= 80)
 
-  // CHECK_PTX71_SM75: call {{.*}} @llvm.nvvm.wmma.m8n8k128.mma.and.popc.row.col.b1
-  // expected-error-re@+1 {{'__bmma_m8n8k128_mma_and_popc_b1' needs target feature (sm_75{{.*}},(ptx71{{.*}}}}
+  // CHECK_PTX71_SM80: call {{.*}} @llvm.nvvm.wmma.m8n8k128.mma.and.popc.row.col.b1
+  // expected-error-re@+1 {{'__bmma_m8n8k128_mma_and_popc_b1' needs target feature (sm_80{{.*}},(ptx71{{.*}}}}
   __bmma_m8n8k128_mma_and_popc_b1(dst, src, src, src, 1);
-#endif // (PTX >= 71) && (SM >= 75) 
+#endif // (PTX >= 71) && (SM >= 80)
 }

clang/test/CodeGen/builtins-nvptx-mma.py

name = "{prefix}_{geom}_mma{b1op}_{suffix}".format(
       b1op = op.b1op.replace(".","_"),
       suffix = suffix)
   return name
 
 def get_required_sm(frag, b1op=""):
   if frag.ptx_type in ["f64", "bf16", "tf32"]:
     return 80
   if frag.ptx_type in ["u4", "s4", "b1"]:
-    if b1op == "_and_popc":
+    if b1op == ".and.popc":
       return 80
     return 75
   if frag.ptx_type in ["s8", "u8"]:
     return 72
   if frag.ptx_type == "s32":
     if frag.geom in ["m8n8k32", "m8n8k128"]: # s4/u4/b1
       return 75
     else:                       # s8/u8
__device__ void test_wmma_buitins(int *src, int *dst,
                                   float *fsrc, float *fdst,
                                   double *dsrc, double *ddst, int ldm) {
 """);
 
   for (ptx, sm), tests in sorted(results.items()):
     print()
     print("#if (PTX >= %d) && (SM >= %d)" % (ptx, sm))
     print(tests)
-    print("#endif // (PTX >= %d) && (SM >= %d) "% (ptx, sm))
+    print("#endif // (PTX >= %d) && (SM >= %d)"% (ptx, sm))
 
   print("}")
 
 parser = argparse.ArgumentParser()
 parser.add_argument("--ptx", type=int, default=60)
 parser.add_argument("--gpu-arch", type=int, default=70)
 args = parser.parse_args()
 ptx_version = args.ptx
 gpu_arch = args.gpu_arch
 
 gen_tests()