clang: Handle MatrixType in hasFloatingRepresentation
authorMatt Arsenault <Matthew.Arsenault@amd.com>
Tue, 28 Feb 2023 11:47:23 +0000 (07:47 -0400)
committerMatt Arsenault <arsenm2@gmail.com>
Wed, 15 Mar 2023 05:13:08 +0000 (01:13 -0400)
Allows applying nofpclass to matrix arguments.

clang/lib/AST/Type.cpp
clang/test/CodeGen/matrix-type-operators-fast-math.c
clang/test/CodeGen/nofpclass.c

index 96e6119..71a098c 100644 (file)
@@ -2154,8 +2154,9 @@ bool Type::isFloatingType() const {
 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();
+  return isFloatingType();
 }
 
 bool Type::isRealFloatingType() const {
index f71faa6..724b282 100644 (file)
@@ -8,7 +8,7 @@ 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]]
@@ -18,7 +18,7 @@ void add_matrix_matrix_double(dx5x5_t a, dx5x5_t b, dx5x5_t 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]]
@@ -28,7 +28,7 @@ void add_compound_assign_matrix_double(dx5x5_t a, dx5x5_t 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]]
@@ -38,7 +38,7 @@ void subtract_compound_assign_matrix_double(dx5x5_t a, dx5x5_t 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
@@ -51,7 +51,7 @@ void add_matrix_scalar_double_float(dx5x5_t a, float 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
@@ -64,7 +64,7 @@ void add_compound_matrix_scalar_double_float(dx5x5_t a, float 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
index 5f28ef8..7fea819 100644 (file)
@@ -1396,40 +1396,40 @@ __m256d call_m256d(__m256d 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) {