clang: Handle MatrixType in hasFloatingRepresentation

[arsenm]

Summary

Allows applying nofpclass to matrix arguments.

[fhahn]

LGTM, thanks!

[arsenm]

ffe12e765cb9ba77dc30ccec89fb163234b98541

clang/lib/AST/Type.cpp

bool Type::isFloatingType() const {
   if (const auto *CT = dyn_cast<ComplexType>(CanonicalType))
     return CT->getElementType()->isFloatingType();
   return false;
 }
 
 bool Type::hasFloatingRepresentation() const {
   if (const auto *VT = dyn_cast<VectorType>(CanonicalType))
     return VT->getElementType()->isFloatingType();
-  else
-    return isFloatingType();
+  if (const auto *MT = dyn_cast<MatrixType>(CanonicalType))
+    return MT->getElementType()->isFloatingType();
 }
 
 bool Type::isRealFloatingType() const {
   if (const auto *BT = dyn_cast<BuiltinType>(CanonicalType))
     return BT->isFloatingPoint();
   return false;
 }
 

clang/test/CodeGen/matrix-type-operators-fast-math.c

 // RUN: %clang_cc1 -ffast-math -fenable-matrix -triple x86_64-apple-darwin %s -emit-llvm -disable-llvm-passes -o - | FileCheck %s
 
 typedef double dx5x5_t __attribute__((matrix_type(5, 5)));
 typedef float fx2x3_t __attribute__((matrix_type(2, 3)));
 typedef int ix9x3_t __attribute__((matrix_type(9, 3)));
 typedef unsigned long long ullx4x2_t __attribute__((matrix_type(4, 2)));
 
 // Floating point matrix/scalar additions.
 
 void add_matrix_matrix_double(dx5x5_t a, dx5x5_t b, dx5x5_t c) {
-  // CHECK-LABEL: define{{.*}} void @add_matrix_matrix_double(<25 x double> noundef %a, <25 x double> noundef %b, <25 x double> noundef %c)
+  // CHECK-LABEL: define{{.*}} void @add_matrix_matrix_double(<25 x double> noundef nofpclass(nan inf) %a, <25 x double> noundef nofpclass(nan inf) %b, <25 x double> noundef nofpclass(nan inf) %c)
   // CHECK:       [[B:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[C:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[RES:%.*]] = fadd reassoc nnan ninf nsz arcp afn <25 x double> [[B]], [[C]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 
   a = b + c;
 }
 
 void add_compound_assign_matrix_double(dx5x5_t a, dx5x5_t b) {
-  // CHECK-LABEL: define{{.*}} void @add_compound_assign_matrix_double(<25 x double> noundef %a, <25 x double> noundef %b)
+  // CHECK-LABEL: define{{.*}} void @add_compound_assign_matrix_double(<25 x double> noundef nofpclass(nan inf) %a, <25 x double> noundef nofpclass(nan inf) %b)
   // CHECK:       [[B:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[A:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[RES:%.*]] = fadd reassoc nnan ninf nsz arcp afn <25 x double> [[A]], [[B]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 
   a += b;
 }
 
 void subtract_compound_assign_matrix_double(dx5x5_t a, dx5x5_t b) {
-  // CHECK-LABEL: define{{.*}} void @subtract_compound_assign_matrix_double(<25 x double> noundef %a, <25 x double> noundef %b)
+  // CHECK-LABEL: define{{.*}} void @subtract_compound_assign_matrix_double(<25 x double> noundef nofpclass(nan inf) %a, <25 x double> noundef nofpclass(nan inf) %b)
   // CHECK:       [[B:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[A:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[RES:%.*]] = fsub reassoc nnan ninf nsz arcp afn <25 x double> [[A]], [[B]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 
   a -= b;
 }
 
 void add_matrix_scalar_double_float(dx5x5_t a, float vf) {
-  // CHECK-LABEL: define{{.*}} void @add_matrix_scalar_double_float(<25 x double> noundef %a, float noundef nofpclass(nan inf) %vf)
+  // CHECK-LABEL: define{{.*}} void @add_matrix_scalar_double_float(<25 x double> noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %vf)
   // CHECK:       [[MATRIX:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[SCALAR:%.*]] = load float, ptr %vf.addr, align 4
   // CHECK-NEXT:  [[SCALAR_EXT:%.*]] = fpext float [[SCALAR]] to double
   // CHECK-NEXT:  [[SCALAR_EMBED:%.*]] = insertelement <25 x double> poison, double [[SCALAR_EXT]], i64 0
   // CHECK-NEXT:  [[SCALAR_EMBED1:%.*]] = shufflevector <25 x double> [[SCALAR_EMBED]], <25 x double> poison, <25 x i32> zeroinitializer
   // CHECK-NEXT:  [[RES:%.*]] = fadd reassoc nnan ninf nsz arcp afn <25 x double> [[MATRIX]], [[SCALAR_EMBED1]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 
   a = a + vf;
 }
 
 void add_compound_matrix_scalar_double_float(dx5x5_t a, float vf) {
-  // CHECK-LABEL: define{{.*}} void @add_compound_matrix_scalar_double_float(<25 x double> noundef %a, float noundef nofpclass(nan inf) %vf)
+  // CHECK-LABEL: define{{.*}} void @add_compound_matrix_scalar_double_float(<25 x double> noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %vf)
   // CHECK:  [[SCALAR:%.*]] = load float, ptr %vf.addr, align 4
   // CHECK-NEXT:  [[SCALAR_EXT:%.*]] = fpext float [[SCALAR]] to double
   // CHECK-NEXT:  [[MATRIX:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[SCALAR_EMBED:%.*]] = insertelement <25 x double> poison, double [[SCALAR_EXT]], i64 0
   // CHECK-NEXT:  [[SCALAR_EMBED1:%.*]] = shufflevector <25 x double> [[SCALAR_EMBED]], <25 x double> poison, <25 x i32> zeroinitializer
   // CHECK-NEXT:  [[RES:%.*]] = fadd reassoc nnan ninf nsz arcp afn <25 x double> [[MATRIX]], [[SCALAR_EMBED1]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 
   a += vf;
 }
 
 void subtract_compound_matrix_scalar_double_float(dx5x5_t a, float vf) {
-  // CHECK-LABEL: define{{.*}} void @subtract_compound_matrix_scalar_double_float(<25 x double> noundef %a, float noundef nofpclass(nan inf) %vf)
+  // CHECK-LABEL: define{{.*}} void @subtract_compound_matrix_scalar_double_float(<25 x double> noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %vf)
   // CHECK:  [[SCALAR:%.*]] = load float, ptr %vf.addr, align 4
   // CHECK-NEXT:  [[SCALAR_EXT:%.*]] = fpext float [[SCALAR]] to double
   // CHECK-NEXT:  [[MATRIX:%.*]] = load <25 x double>, ptr {{.*}}, align 8
   // CHECK-NEXT:  [[SCALAR_EMBED:%.*]] = insertelement <25 x double> poison, double [[SCALAR_EXT]], i64 0
   // CHECK-NEXT:  [[SCALAR_EMBED1:%.*]] = shufflevector <25 x double> [[SCALAR_EMBED]], <25 x double> poison, <25 x i32> zeroinitializer
   // CHECK-NEXT:  [[RES:%.*]] = fsub reassoc nnan ninf nsz arcp afn <25 x double> [[MATRIX]], [[SCALAR_EMBED1]]
   // CHECK-NEXT:  store <25 x double> [[RES]], ptr {{.*}}, align 8
 

clang/test/CodeGen/nofpclass.c

 // NOINFS-NEXT:    [[CALL:%.*]] = call ninf nofpclass(inf) <4 x double> (<4 x double>, ...) @extern_m256d(<4 x double> noundef nofpclass(inf) [[TMP0]], <4 x double> noundef nofpclass(inf) [[TMP1]])
 // NOINFS-NEXT:    ret <4 x double> [[CALL]]
 //
 __m256d call_m256d(__m256d x) {
   return extern_m256d(x, x);
 }
 
 // CFINITEONLY: Function Attrs: noinline nounwind optnone
-// CFINITEONLY-LABEL: define dso_local <25 x double> @call_matrix
-// CFINITEONLY-SAME: (<25 x double> noundef [[X:%.*]]) #[[ATTR6:[0-9]+]] {
+// CFINITEONLY-LABEL: define dso_local nofpclass(nan inf) <25 x double> @call_matrix
+// CFINITEONLY-SAME: (<25 x double> noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR6:[0-9]+]] {
 // CFINITEONLY-NEXT:  entry:
 // CFINITEONLY-NEXT:    [[X_ADDR:%.*]] = alloca [25 x double], align 8
 // CFINITEONLY-NEXT:    store <25 x double> [[X]], ptr [[X_ADDR]], align 8
 // CFINITEONLY-NEXT:    [[TMP0:%.*]] = load <25 x double>, ptr [[X_ADDR]], align 8
-// CFINITEONLY-NEXT:    [[CALL:%.*]] = call nnan ninf <25 x double> @extern_matrix(<25 x double> noundef [[TMP0]])
+// CFINITEONLY-NEXT:    [[CALL:%.*]] = call nnan ninf nofpclass(nan inf) <25 x double> @extern_matrix(<25 x double> noundef nofpclass(nan inf) [[TMP0]])
 // CFINITEONLY-NEXT:    ret <25 x double> [[CALL]]
 //
 // CLFINITEONLY: Function Attrs: convergent norecurse nounwind
-// CLFINITEONLY-LABEL: define dso_local <25 x double> @call_matrix
-// CLFINITEONLY-SAME: (<25 x double> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR9:[0-9]+]] {
+// CLFINITEONLY-LABEL: define dso_local nofpclass(nan inf) <25 x double> @call_matrix
+// CLFINITEONLY-SAME: (<25 x double> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR9:[0-9]+]] {
 // CLFINITEONLY-NEXT:  entry:
-// CLFINITEONLY-NEXT:    [[CALL:%.*]] = tail call nnan ninf <25 x double> @extern_matrix(<25 x double> noundef [[X]]) #[[ATTR10]]
+// CLFINITEONLY-NEXT:    [[CALL:%.*]] = tail call nnan ninf nofpclass(nan inf) <25 x double> @extern_matrix(<25 x double> noundef nofpclass(nan inf) [[X]]) #[[ATTR10]]
 // CLFINITEONLY-NEXT:    ret <25 x double> [[CALL]]
 //
 // NONANS: Function Attrs: noinline nounwind optnone
-// NONANS-LABEL: define dso_local <25 x double> @call_matrix
-// NONANS-SAME: (<25 x double> noundef [[X:%.*]]) #[[ATTR6:[0-9]+]] {
+// NONANS-LABEL: define dso_local nofpclass(nan) <25 x double> @call_matrix
+// NONANS-SAME: (<25 x double> noundef nofpclass(nan) [[X:%.*]]) #[[ATTR6:[0-9]+]] {
 // NONANS-NEXT:  entry:
 // NONANS-NEXT:    [[X_ADDR:%.*]] = alloca [25 x double], align 8
 // NONANS-NEXT:    store <25 x double> [[X]], ptr [[X_ADDR]], align 8
 // NONANS-NEXT:    [[TMP0:%.*]] = load <25 x double>, ptr [[X_ADDR]], align 8
-// NONANS-NEXT:    [[CALL:%.*]] = call nnan <25 x double> @extern_matrix(<25 x double> noundef [[TMP0]])
+// NONANS-NEXT:    [[CALL:%.*]] = call nnan nofpclass(nan) <25 x double> @extern_matrix(<25 x double> noundef nofpclass(nan) [[TMP0]])
 // NONANS-NEXT:    ret <25 x double> [[CALL]]
 //
 // NOINFS: Function Attrs: noinline nounwind optnone
-// NOINFS-LABEL: define dso_local <25 x double> @call_matrix
-// NOINFS-SAME: (<25 x double> noundef [[X:%.*]]) #[[ATTR6:[0-9]+]] {
+// NOINFS-LABEL: define dso_local nofpclass(inf) <25 x double> @call_matrix
+// NOINFS-SAME: (<25 x double> noundef nofpclass(inf) [[X:%.*]]) #[[ATTR6:[0-9]+]] {
 // NOINFS-NEXT:  entry:
 // NOINFS-NEXT:    [[X_ADDR:%.*]] = alloca [25 x double], align 8
 // NOINFS-NEXT:    store <25 x double> [[X]], ptr [[X_ADDR]], align 8
 // NOINFS-NEXT:    [[TMP0:%.*]] = load <25 x double>, ptr [[X_ADDR]], align 8
-// NOINFS-NEXT:    [[CALL:%.*]] = call ninf <25 x double> @extern_matrix(<25 x double> noundef [[TMP0]])
+// NOINFS-NEXT:    [[CALL:%.*]] = call ninf nofpclass(inf) <25 x double> @extern_matrix(<25 x double> noundef nofpclass(inf) [[TMP0]])
 // NOINFS-NEXT:    ret <25 x double> [[CALL]]
 //
 dx5x5_t call_matrix(dx5x5_t x) {
   return extern_matrix(x);
 }