yaxunl updated this revision to Diff 526258.
yaxunl marked 2 inline comments as done.
yaxunl added a comment.
revised by Artem's comments
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D150985/new/
https://reviews.llvm.org/D150985
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/CodeGen/CGAtomic.cpp
clang/lib/Sema/SemaChecking.cpp
clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
clang/test/Sema/atomic-ops.c
clang/test/SemaOpenCL/atomic-ops.cl
Index: clang/test/SemaOpenCL/atomic-ops.cl
===================================================================
--- clang/test/SemaOpenCL/atomic-ops.cl
+++ clang/test/SemaOpenCL/atomic-ops.cl
@@ -61,8 +61,10 @@
__opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_fetch_min(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
- __opencl_atomic_fetch_max(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+ __opencl_atomic_fetch_min(f, 1, memory_order_seq_cst, memory_scope_work_group);
+ __opencl_atomic_fetch_max(f, 1, memory_order_seq_cst, memory_scope_work_group);
+ __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group);
+ __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group);
bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
Index: clang/test/Sema/atomic-ops.c
===================================================================
--- clang/test/Sema/atomic-ops.c
+++ clang/test/Sema/atomic-ops.c
@@ -199,14 +199,24 @@
__c11_atomic_fetch_add(f, 1.0f, memory_order_seq_cst);
__c11_atomic_fetch_add(d, 1.0, memory_order_seq_cst);
__c11_atomic_fetch_add(ld, 1.0, memory_order_seq_cst); // fp80-error {{must be a pointer to atomic integer, pointer or supported floating point type}}
+ __c11_atomic_fetch_min(i, 1, memory_order_seq_cst);
+ __c11_atomic_fetch_min(p, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer or supported floating point type}}
+ __c11_atomic_fetch_min(f, 1.0f, memory_order_seq_cst);
+ __c11_atomic_fetch_min(d, 1.0, memory_order_seq_cst);
+ __c11_atomic_fetch_min(ld, 1.0, memory_order_seq_cst); // fp80-error {{must be a pointer to atomic integer or supported floating point type}}
+ __c11_atomic_fetch_max(i, 1, memory_order_seq_cst);
+ __c11_atomic_fetch_max(p, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer or supported floating point type}}
+ __c11_atomic_fetch_max(f, 1.0f, memory_order_seq_cst);
+ __c11_atomic_fetch_max(d, 1.0, memory_order_seq_cst);
+ __c11_atomic_fetch_max(ld, 1.0, memory_order_seq_cst); // fp80-error {{must be a pointer to atomic integer or supported floating point type}}
__atomic_fetch_add(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer, pointer or supported floating point type}}
__atomic_fetch_sub(I, 3, memory_order_seq_cst);
__atomic_fetch_sub(P, 3, memory_order_seq_cst);
__atomic_fetch_sub(D, 3, memory_order_seq_cst);
__atomic_fetch_sub(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
- __atomic_fetch_min(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
- __atomic_fetch_max(P, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
+ __atomic_fetch_min(D, 3, memory_order_seq_cst);
+ __atomic_fetch_max(P, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer or supported floating point type}}
__atomic_fetch_max(p, 3); // expected-error {{too few arguments to function call, expected 3, have 2}}
__c11_atomic_fetch_and(i, 1, memory_order_seq_cst);
Index: clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
+++ clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -1,29 +1,98 @@
-// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
// RUN: -fnative-half-arguments-and-returns | FileCheck %s
+// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -target-cpu gfx1100 -fnative-half-type \
+// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefix=SAFE %s
+
+// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -target-cpu gfx940 -fnative-half-type \
+// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \
+// RUN: | FileCheck -check-prefix=UNSAFE %s
+
// REQUIRES: amdgpu-registered-target
#include "Inputs/cuda.h"
#include <stdatomic.h>
-__device__ float ffp1(float *p) {
+__global__ void ffp1(float *p) {
// CHECK-LABEL: @_Z4ffp1Pf
// CHECK: atomicrmw fadd ptr {{.*}} monotonic
- return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+ // CHECK: atomicrmw fmax ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic
+ // SAFE: _Z4ffp1Pf
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // UNSAFE: _Z4ffp1Pf
+ // UNSAFE: global_atomic_add_f32
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+ __atomic_fetch_max(p, 1.0f, memory_order_relaxed);
+ __atomic_fetch_min(p, 1.0f, memory_order_relaxed);
+ __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
}
-__device__ double ffp2(double *p) {
+__global__ void ffp2(double *p) {
// CHECK-LABEL: @_Z4ffp2Pd
// CHECK: atomicrmw fsub ptr {{.*}} monotonic
- return __atomic_fetch_sub(p, 1.0, memory_order_relaxed);
+ // CHECK: atomicrmw fmax ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic
+ // SAFE: _Z4ffp2Pd
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // UNSAFE: _Z4ffp2Pd
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ __atomic_fetch_sub(p, 1.0, memory_order_relaxed);
+ __atomic_fetch_max(p, 1.0, memory_order_relaxed);
+ __atomic_fetch_min(p, 1.0, memory_order_relaxed);
+ __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
}
// long double is the same as double for amdgcn.
-__device__ long double ffp3(long double *p) {
+__global__ void ffp3(long double *p) {
// CHECK-LABEL: @_Z4ffp3Pe
// CHECK: atomicrmw fsub ptr {{.*}} monotonic
- return __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);
+ // CHECK: atomicrmw fmax ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic
+ // SAFE: _Z4ffp3Pe
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // SAFE: global_atomic_cmpswap_b64
+ // UNSAFE: _Z4ffp3Pe
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ // UNSAFE: global_atomic_cmpswap_x2
+ __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);
+ __atomic_fetch_max(p, 1.0L, memory_order_relaxed);
+ __atomic_fetch_min(p, 1.0L, memory_order_relaxed);
+ __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
}
__device__ double ffp4(double *p, float f) {
@@ -39,3 +108,29 @@
// CHECK: atomicrmw fsub ptr {{.*}} monotonic
return __atomic_fetch_sub(p, i, memory_order_relaxed);
}
+
+__global__ void ffp6(_Float16 *p) {
+ // CHECK-LABEL: @_Z4ffp6PDF16
+ // CHECK: atomicrmw fadd ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmax ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} monotonic
+ // CHECK: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic
+ // CHECK: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic
+ // SAFE: _Z4ffp6PDF16
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // UNSAFE: _Z4ffp6PDF16
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ __atomic_fetch_add(p, 1.0, memory_order_relaxed);
+ __atomic_fetch_max(p, 1.0, memory_order_relaxed);
+ __atomic_fetch_min(p, 1.0, memory_order_relaxed);
+ __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
+}
Index: clang/lib/Sema/SemaChecking.cpp
===================================================================
--- clang/lib/Sema/SemaChecking.cpp
+++ clang/lib/Sema/SemaChecking.cpp
@@ -6400,7 +6400,16 @@
Op == AtomicExpr::AO__atomic_store_n ||
Op == AtomicExpr::AO__atomic_exchange_n ||
Op == AtomicExpr::AO__atomic_compare_exchange_n;
- bool IsAddSub = false;
+ // Bit mask for allowed value types for atomic arithmetic operations.
+ // Add/sub allow integer, pointer and floating point.
+ // Min/max allow integer and floating point.
+ // Other arithmetic operations allow integer.
+ enum ArithOpAllowedValueType {
+ AOAVT_Integer = 1,
+ AOAVT_Pointer = 2,
+ AOAVT_FP = 4,
+ };
+ unsigned ArithAllows = AOAVT_Integer;
switch (Op) {
case AtomicExpr::AO__c11_atomic_init:
@@ -6426,18 +6435,29 @@
case AtomicExpr::AO__atomic_store_n:
Form = Copy;
break;
- case AtomicExpr::AO__hip_atomic_fetch_add:
- case AtomicExpr::AO__hip_atomic_fetch_min:
- case AtomicExpr::AO__hip_atomic_fetch_max:
- case AtomicExpr::AO__c11_atomic_fetch_add:
- case AtomicExpr::AO__c11_atomic_fetch_sub:
- case AtomicExpr::AO__opencl_atomic_fetch_add:
- case AtomicExpr::AO__opencl_atomic_fetch_sub:
case AtomicExpr::AO__atomic_fetch_add:
case AtomicExpr::AO__atomic_fetch_sub:
case AtomicExpr::AO__atomic_add_fetch:
case AtomicExpr::AO__atomic_sub_fetch:
- IsAddSub = true;
+ case AtomicExpr::AO__c11_atomic_fetch_add:
+ case AtomicExpr::AO__c11_atomic_fetch_sub:
+ case AtomicExpr::AO__opencl_atomic_fetch_add:
+ case AtomicExpr::AO__opencl_atomic_fetch_sub:
+ case AtomicExpr::AO__hip_atomic_fetch_add:
+ ArithAllows = AOAVT_Integer | AOAVT_Pointer | AOAVT_FP;
+ Form = Arithmetic;
+ break;
+ case AtomicExpr::AO__atomic_fetch_max:
+ case AtomicExpr::AO__atomic_fetch_min:
+ case AtomicExpr::AO__atomic_max_fetch:
+ case AtomicExpr::AO__atomic_min_fetch:
+ case AtomicExpr::AO__c11_atomic_fetch_max:
+ case AtomicExpr::AO__c11_atomic_fetch_min:
+ case AtomicExpr::AO__opencl_atomic_fetch_max:
+ case AtomicExpr::AO__opencl_atomic_fetch_min:
+ case AtomicExpr::AO__hip_atomic_fetch_max:
+ case AtomicExpr::AO__hip_atomic_fetch_min:
+ ArithAllows = AOAVT_Integer | AOAVT_FP;
Form = Arithmetic;
break;
case AtomicExpr::AO__c11_atomic_fetch_and:
@@ -6460,16 +6480,6 @@
case AtomicExpr::AO__atomic_nand_fetch:
Form = Arithmetic;
break;
- case AtomicExpr::AO__c11_atomic_fetch_min:
- case AtomicExpr::AO__c11_atomic_fetch_max:
- case AtomicExpr::AO__opencl_atomic_fetch_min:
- case AtomicExpr::AO__opencl_atomic_fetch_max:
- case AtomicExpr::AO__atomic_min_fetch:
- case AtomicExpr::AO__atomic_max_fetch:
- case AtomicExpr::AO__atomic_fetch_min:
- case AtomicExpr::AO__atomic_fetch_max:
- Form = Arithmetic;
- break;
case AtomicExpr::AO__c11_atomic_exchange:
case AtomicExpr::AO__hip_atomic_exchange:
@@ -6557,13 +6567,16 @@
if (Form == Arithmetic) {
// GCC does not enforce these rules for GNU atomics, but we do to help catch
// trivial type errors.
- auto IsAllowedValueType = [&](QualType ValType) {
+ auto IsAllowedValueType = [&](QualType ValType,
+ unsigned AllowedType) -> bool {
if (ValType->isIntegerType())
- return true;
+ return AllowedType & AOAVT_Integer;
if (ValType->isPointerType())
- return true;
+ return AllowedType & AOAVT_Pointer;
if (!ValType->isFloatingType())
return false;
+ if (!(AllowedType & AOAVT_FP))
+ return false;
// LLVM Parser does not allow atomicrmw with x86_fp80 type.
if (ValType->isSpecificBuiltinType(BuiltinType::LongDouble) &&
&Context.getTargetInfo().getLongDoubleFormat() ==
@@ -6571,13 +6584,14 @@
return false;
return true;
};
- if (IsAddSub && !IsAllowedValueType(ValType)) {
- Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_ptr_or_fp)
- << IsC11 << Ptr->getType() << Ptr->getSourceRange();
- return ExprError();
- }
- if (!IsAddSub && !ValType->isIntegerType()) {
- Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int)
+ if (!IsAllowedValueType(ValType, ArithAllows)) {
+ assert(ArithAllows & AOAVT_Integer);
+ auto DID = ArithAllows & AOAVT_FP
+ ? (ArithAllows & AOAVT_Pointer
+ ? diag::err_atomic_op_needs_atomic_int_ptr_or_fp
+ : diag::err_atomic_op_needs_atomic_int_or_fp)
+ : diag::err_atomic_op_needs_atomic_int;
+ Diag(ExprRange.getBegin(), DID)
<< IsC11 << Ptr->getType() << Ptr->getSourceRange();
return ExprError();
}
Index: clang/lib/CodeGen/CGAtomic.cpp
===================================================================
--- clang/lib/CodeGen/CGAtomic.cpp
+++ clang/lib/CodeGen/CGAtomic.cpp
@@ -636,8 +636,11 @@
case AtomicExpr::AO__hip_atomic_fetch_min:
case AtomicExpr::AO__opencl_atomic_fetch_min:
case AtomicExpr::AO__atomic_fetch_min:
- Op = E->getValueType()->isSignedIntegerType() ? llvm::AtomicRMWInst::Min
- : llvm::AtomicRMWInst::UMin;
+ Op = E->getValueType()->isFloatingType()
+ ? llvm::AtomicRMWInst::FMin
+ : (E->getValueType()->isSignedIntegerType()
+ ? llvm::AtomicRMWInst::Min
+ : llvm::AtomicRMWInst::UMin);
break;
case AtomicExpr::AO__atomic_max_fetch:
@@ -647,8 +650,11 @@
case AtomicExpr::AO__hip_atomic_fetch_max:
case AtomicExpr::AO__opencl_atomic_fetch_max:
case AtomicExpr::AO__atomic_fetch_max:
- Op = E->getValueType()->isSignedIntegerType() ? llvm::AtomicRMWInst::Max
- : llvm::AtomicRMWInst::UMax;
+ Op = E->getValueType()->isFloatingType()
+ ? llvm::AtomicRMWInst::FMax
+ : (E->getValueType()->isSignedIntegerType()
+ ? llvm::AtomicRMWInst::Max
+ : llvm::AtomicRMWInst::UMax);
break;
case AtomicExpr::AO__atomic_and_fetch:
@@ -916,9 +922,19 @@
}
[[fallthrough]];
case AtomicExpr::AO__atomic_fetch_add:
+ case AtomicExpr::AO__atomic_fetch_max:
+ case AtomicExpr::AO__atomic_fetch_min:
case AtomicExpr::AO__atomic_fetch_sub:
case AtomicExpr::AO__atomic_add_fetch:
+ case AtomicExpr::AO__atomic_max_fetch:
+ case AtomicExpr::AO__atomic_min_fetch:
case AtomicExpr::AO__atomic_sub_fetch:
+ case AtomicExpr::AO__c11_atomic_fetch_max:
+ case AtomicExpr::AO__c11_atomic_fetch_min:
+ case AtomicExpr::AO__opencl_atomic_fetch_max:
+ case AtomicExpr::AO__opencl_atomic_fetch_min:
+ case AtomicExpr::AO__hip_atomic_fetch_max:
+ case AtomicExpr::AO__hip_atomic_fetch_min:
ShouldCastToIntPtrTy = !MemTy->isFloatingType();
[[fallthrough]];
@@ -934,13 +950,9 @@
case AtomicExpr::AO__c11_atomic_fetch_or:
case AtomicExpr::AO__c11_atomic_fetch_xor:
case AtomicExpr::AO__c11_atomic_fetch_nand:
- case AtomicExpr::AO__c11_atomic_fetch_max:
- case AtomicExpr::AO__c11_atomic_fetch_min:
case AtomicExpr::AO__opencl_atomic_fetch_and:
case AtomicExpr::AO__opencl_atomic_fetch_or:
case AtomicExpr::AO__opencl_atomic_fetch_xor:
- case AtomicExpr::AO__opencl_atomic_fetch_min:
- case AtomicExpr::AO__opencl_atomic_fetch_max:
case AtomicExpr::AO__atomic_fetch_and:
case AtomicExpr::AO__hip_atomic_fetch_and:
case AtomicExpr::AO__atomic_fetch_or:
@@ -952,12 +964,6 @@
case AtomicExpr::AO__atomic_or_fetch:
case AtomicExpr::AO__atomic_xor_fetch:
case AtomicExpr::AO__atomic_nand_fetch:
- case AtomicExpr::AO__atomic_max_fetch:
- case AtomicExpr::AO__atomic_min_fetch:
- case AtomicExpr::AO__atomic_fetch_max:
- case AtomicExpr::AO__hip_atomic_fetch_max:
- case AtomicExpr::AO__atomic_fetch_min:
- case AtomicExpr::AO__hip_atomic_fetch_min:
Val1 = EmitValToTemp(*this, E->getVal1());
break;
}
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8523,6 +8523,9 @@
def err_atomic_op_needs_atomic_int_or_ptr : Error<
"address argument to atomic operation must be a pointer to %select{|atomic }0"
"integer or pointer (%1 invalid)">;
+def err_atomic_op_needs_atomic_int_or_fp : Error<
+ "address argument to atomic operation must be a pointer to %select{|atomic }0"
+ "integer or supported floating point type (%1 invalid)">;
def err_atomic_op_needs_atomic_int : Error<
"address argument to atomic operation must be a pointer to "
"%select{|atomic }0integer (%1 invalid)">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
