[clang] [AMDGPU] Allow overload of __builtin_amdgcn_mov/update_dpp (PR #112447)
Stanislav Mekhanoshin via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 17 14:33:17 PDT 2024
https://github.com/rampitec updated https://github.com/llvm/llvm-project/pull/112447
>From 761b3e21748dd3a7b53cd0ead745943213317eb4 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Tue, 15 Oct 2024 15:23:28 -0700
Subject: [PATCH 1/5] [AMDGPU] Allow overload of
__builtin_amdgcn_mov/update_dpp
We need to support 64-bit data types (intrinsics do support it).
We are also silently converting FP to integer argument now,
also fixed.
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 +-
clang/lib/CodeGen/CGBuiltin.cpp | 26 ++++++---
.../test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 54 +++++++++++++++++--
3 files changed, 71 insertions(+), 13 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index c02970f55b22e1..e887213aa945e6 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -224,8 +224,8 @@ TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "WUi", "n", "s-memrealtime")
-TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
-TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp")
+TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nct", "dpp")
+TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nct", "dpp")
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "gfx8-insts")
TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index c563f2618b42c8..22e707ca552d28 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19038,15 +19038,27 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
ASTContext::GetBuiltinTypeError Error;
getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);
assert(Error == ASTContext::GE_None && "Should not codegen an error");
+ llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());
+ llvm::Type *IntTy = llvm::IntegerType::get(
+ Builder.getContext(), DataTy->getPrimitiveSizeInBits());
+ Function *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
+ assert(E->getNumArgs() == 5 || E->getNumArgs() == 6);
+ bool InsertOld = E->getNumArgs() == 5;
+ if (InsertOld)
+ Args.push_back(llvm::PoisonValue::get(IntTy));
for (unsigned I = 0; I != E->getNumArgs(); ++I) {
- Args.push_back(EmitScalarOrConstFoldImmArg(ICEArguments, I, E));
+ llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);
+ llvm::Type *ExpTy =
+ F->getFunctionType()->getFunctionParamType(I + InsertOld);
+ if (V->getType() != ExpTy)
+ V = Builder.CreateTruncOrBitCast(V, ExpTy);
+ Args.push_back(V);
}
- assert(Args.size() == 5 || Args.size() == 6);
- if (Args.size() == 5)
- Args.insert(Args.begin(), llvm::PoisonValue::get(Args[0]->getType()));
- Function *F =
- CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
- return Builder.CreateCall(F, Args);
+ llvm::Value *V = Builder.CreateCall(F, Args);
+ if (!DataTy->isIntegerTy())
+ V = Builder.CreateBitCast(V, DataTy);
+ return V;
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
case AMDGPU::BI__builtin_amdgcn_permlanex16:
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index 5bd8f77a5930c4..a2de7bb953d584 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -102,20 +102,66 @@ void test_s_dcache_wb()
__builtin_amdgcn_s_dcache_wb();
}
-// CHECK-LABEL: @test_mov_dpp
+// CHECK-LABEL: @test_mov_dpp_int
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %src, i32 0, i32 0, i32 0, i1 false)
-void test_mov_dpp(global int* out, int src)
+void test_mov_dpp_int(global int* out, int src)
{
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
}
-// CHECK-LABEL: @test_update_dpp
+// CHECK-LABEL: @test_mov_dpp_long
+// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %x, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: store i64 %0,
+void test_mov_dpp_long(long x, global long *p) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_mov_dpp_float
+// CHECK: %0 = bitcast float %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: store i32 %1,
+void test_mov_dpp_float(float x, global float *p) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_mov_dpp_double
+// CHECK: %0 = bitcast double %x to i64
+// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: store i64 %1,
+void test_mov_dpp_double(double x, global double *p) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_update_dpp_int
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
-void test_update_dpp(global int* out, int arg1, int arg2)
+void test_update_dpp_int(global int* out, int arg1, int arg2)
{
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
}
+// CHECK-LABEL: @test_update_dpp_long
+// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %x, i64 %x, i32 257, i32 15, i32 15, i1 false)
+// CHECk-NEXT: store i64 %0,
+void test_update_dpp_long(long x, global long *p) {
+ *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_update_dpp_float
+// CHECK: %0 = bitcast float %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: store i32 %1,
+void test_update_dpp_float(float x, global float *p) {
+ *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_update_dpp_double
+// CHECK: %0 = bitcast double %x to i64
+// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %0, i64 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: store i64 %1,
+void test_update_dpp_double(double x, global double *p) {
+ *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
+}
+
// CHECK-LABEL: @test_ds_fadd
// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
// CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
>From cb9a29e6dfa16ec3d8315ee87470eec0af98ac72 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Tue, 15 Oct 2024 15:40:46 -0700
Subject: [PATCH 2/5] clang-format
---
clang/lib/CodeGen/CGBuiltin.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 22e707ca552d28..ee12fbc66a0e51 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19041,8 +19041,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());
llvm::Type *IntTy = llvm::IntegerType::get(
Builder.getContext(), DataTy->getPrimitiveSizeInBits());
- Function *F =
- CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
+ Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
assert(E->getNumArgs() == 5 || E->getNumArgs() == 6);
bool InsertOld = E->getNumArgs() == 5;
if (InsertOld)
>From 9b07ba98afcce1cd911cdf982b9e59bfa5827f60 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Tue, 15 Oct 2024 15:58:27 -0700
Subject: [PATCH 3/5] Simplify casts, builder will not produce a cast to the
same type.
---
clang/lib/CodeGen/CGBuiltin.cpp | 9 ++-------
1 file changed, 2 insertions(+), 7 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index ee12fbc66a0e51..d2daab4841d759 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19050,14 +19050,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);
llvm::Type *ExpTy =
F->getFunctionType()->getFunctionParamType(I + InsertOld);
- if (V->getType() != ExpTy)
- V = Builder.CreateTruncOrBitCast(V, ExpTy);
- Args.push_back(V);
+ Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy));
}
- llvm::Value *V = Builder.CreateCall(F, Args);
- if (!DataTy->isIntegerTy())
- V = Builder.CreateBitCast(V, DataTy);
- return V;
+ return Builder.CreateBitCast(Builder.CreateCall(F, Args), DataTy);
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
case AMDGPU::BI__builtin_amdgcn_permlanex16:
>From 27ab5cc41a4681d2b767de63fa1d31bd1784df76 Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Wed, 16 Oct 2024 12:35:57 -0700
Subject: [PATCH 4/5] Added sema checks
---
clang/lib/Sema/SemaAMDGPU.cpp | 28 +++++++++++++++
.../SemaOpenCL/builtins-amdgcn-error-gfx9.cl | 34 +++++++++++++++++++
2 files changed, 62 insertions(+)
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index d11bc9eec33019..9126c5ec3cd7ba 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -63,6 +63,34 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
OrderIndex = 0;
ScopeIndex = 1;
break;
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
+ if (SemaRef.checkArgCountRange(TheCall, 5, 5))
+ return true;
+ Expr *ValArg = TheCall->getArg(0);
+ QualType Ty = ValArg->getType();
+ if (!Ty->isArithmeticType()) {
+ SemaRef.Diag(ValArg->getBeginLoc(),
+ diag::err_typecheck_cond_expect_int_float)
+ << Ty << ValArg->getSourceRange();
+ return true;
+ }
+ return false;
+ }
+ case AMDGPU::BI__builtin_amdgcn_update_dpp: {
+ if (SemaRef.checkArgCountRange(TheCall, 6, 6))
+ return true;
+ for (unsigned I = 0; I != 2; ++I) {
+ Expr *ValArg = TheCall->getArg(I);
+ QualType Ty = ValArg->getType();
+ if (!Ty->isArithmeticType()) {
+ SemaRef.Diag(ValArg->getBeginLoc(),
+ diag::err_typecheck_cond_expect_int_float)
+ << Ty << ValArg->getSourceRange();
+ return true;
+ }
+ }
+ return false;
+ }
default:
return false;
}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
index c9fd8ab2cae86c..8dabcabbad0df1 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
@@ -7,3 +7,37 @@ void test_gfx9_fmed3h(global half *out, half a, half b, half c)
{
*out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
}
+
+void test_mov_dpp(global int* out, int src, int i)
+{
+ *out = __builtin_amdgcn_mov_dpp(src, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, i, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, i, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, i); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0.1, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0.1, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0.1, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, 0.1); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0); // expected-error{{too few arguments to function call, expected 5, have 4}}
+ *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false, 1); // expected-error{{too many arguments to function call, expected at most 5, have 6}}
+ *out = __builtin_amdgcn_mov_dpp(out, 0, 0, 0, false); // expected-error{{used type '__global int *__private' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_mov_dpp("aa", 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
+}
+
+void test_update_dpp(global int* out, int arg1, int arg2, int i)
+{
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, i, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, i, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, i); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0.1, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0.1, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0.1, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, 0.1); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0); // expected-error{{too few arguments to function call, expected 6, have 5}}
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false, 1); // expected-error{{too many arguments to function call, expected at most 6, have 7}}
+ *out = __builtin_amdgcn_update_dpp(out, arg2, 0, 0, 0, false); // expected-error{{used type '__global int *__private' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(arg1, out, 0, 0, 0, false); // expected-error{{used type '__global int *__private' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp("aa", arg2, 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(arg1, "aa", 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
+}
>From 9cc15d55ca6faa8c5460f6ed4de76ba7e6fd427d Mon Sep 17 00:00:00 2001
From: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
Date: Thu, 17 Oct 2024 14:32:38 -0700
Subject: [PATCH 5/5] Fixed subdword and complex handling, added tests
---
clang/lib/CodeGen/CGBuiltin.cpp | 15 ++++-
clang/lib/Sema/SemaAMDGPU.cpp | 6 +-
.../test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 56 +++++++++++++++++++
.../SemaOpenCL/builtins-amdgcn-error-gfx9.cl | 19 ++++++-
4 files changed, 90 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 9db46579d6d7c8..5831352ea20b33 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19043,8 +19043,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);
assert(Error == ASTContext::GE_None && "Should not codegen an error");
llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());
+ unsigned Size = DataTy->getPrimitiveSizeInBits();
llvm::Type *IntTy = llvm::IntegerType::get(
- Builder.getContext(), DataTy->getPrimitiveSizeInBits());
+ Builder.getContext(), std::max(Size, 32u));
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
assert(E->getNumArgs() == 5 || E->getNumArgs() == 6);
bool InsertOld = E->getNumArgs() == 5;
@@ -19052,11 +19053,21 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
Args.push_back(llvm::PoisonValue::get(IntTy));
for (unsigned I = 0; I != E->getNumArgs(); ++I) {
llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);
+ if (I <= !InsertOld && Size < 32) {
+ if (!DataTy->isIntegerTy())
+ V = Builder.CreateBitCast(
+ V, llvm::IntegerType::get(Builder.getContext(), Size));
+ V = Builder.CreateZExtOrBitCast(V, IntTy);
+ }
llvm::Type *ExpTy =
F->getFunctionType()->getFunctionParamType(I + InsertOld);
Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy));
}
- return Builder.CreateBitCast(Builder.CreateCall(F, Args), DataTy);
+ Value *V = Builder.CreateCall(F, Args);
+ if (Size < 32 && !DataTy->isIntegerTy())
+ V = Builder.CreateTrunc(
+ V, llvm::IntegerType::get(Builder.getContext(), Size));
+ return Builder.CreateTruncOrBitCast(V, DataTy);
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
case AMDGPU::BI__builtin_amdgcn_permlanex16:
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 9126c5ec3cd7ba..dd0fdf6f1dc38f 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -68,7 +68,8 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
return true;
Expr *ValArg = TheCall->getArg(0);
QualType Ty = ValArg->getType();
- if (!Ty->isArithmeticType()) {
+ // TODO: Vectors can also be supported.
+ if (!Ty->isArithmeticType() || Ty->isAnyComplexType()) {
SemaRef.Diag(ValArg->getBeginLoc(),
diag::err_typecheck_cond_expect_int_float)
<< Ty << ValArg->getSourceRange();
@@ -82,7 +83,8 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
for (unsigned I = 0; I != 2; ++I) {
Expr *ValArg = TheCall->getArg(I);
QualType Ty = ValArg->getType();
- if (!Ty->isArithmeticType()) {
+ // TODO: Vectors can also be supported.
+ if (!Ty->isArithmeticType() || Ty->isAnyComplexType()) {
SemaRef.Diag(ValArg->getBeginLoc(),
diag::err_typecheck_cond_expect_int_float)
<< Ty << ValArg->getSourceRange();
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index a2de7bb953d584..65b54c1d552742 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -132,6 +132,34 @@ void test_mov_dpp_double(double x, global double *p) {
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
}
+// CHECK-LABEL: @test_mov_dpp_short
+// CHECK: %0 = zext i16 %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %2 = trunc i32 %1 to i16
+// CHECK-NEXT: store i16 %2,
+void test_mov_dpp_short(short x, global short *p) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_mov_dpp_char
+// CHECK: %0 = zext i8 %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %2 = trunc i32 %1 to i8
+// CHECK-NEXT: store i8 %2,
+void test_mov_dpp_char(char x, global char *p) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_mov_dpp_half
+// CHECK: %0 = load i16,
+// CHECK: %1 = zext i16 %0 to i32
+// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %1, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %3 = trunc i32 %2 to i16
+// CHECK-NEXT: store i16 %3,
+void test_mov_dpp_half(half *x, global half *p) {
+ *p = __builtin_amdgcn_mov_dpp(*x, 0x101, 0xf, 0xf, 0);
+}
+
// CHECK-LABEL: @test_update_dpp_int
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
void test_update_dpp_int(global int* out, int arg1, int arg2)
@@ -162,6 +190,34 @@ void test_update_dpp_double(double x, global double *p) {
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
}
+// CHECK-LABEL: @test_update_dpp_short
+// CHECK: %0 = zext i16 %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %2 = trunc i32 %1 to i16
+// CHECK-NEXT: store i16 %2,
+void test_update_dpp_short(short x, global short *p) {
+ *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_update_dpp_char
+// CHECK: %0 = zext i8 %x to i32
+// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %2 = trunc i32 %1 to i8
+// CHECK-NEXT: store i8 %2,
+void test_update_dpp_char(char x, global char *p) {
+ *p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
+}
+
+// CHECK-LABEL: @test_update_dpp_half
+// CHECK: %0 = load i16,
+// CHECK: %1 = zext i16 %0 to i32
+// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %1, i32 %1, i32 257, i32 15, i32 15, i1 false)
+// CHECK-NEXT: %3 = trunc i32 %2 to i16
+// CHECK-NEXT: store i16 %3,
+void test_update_dpp_half(half *x, global half *p) {
+ *p = __builtin_amdgcn_update_dpp(*x, *x, 0x101, 0xf, 0xf, 0);
+}
+
// CHECK-LABEL: @test_ds_fadd
// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
// CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
index 8dabcabbad0df1..6a6a0436169258 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
@@ -3,12 +3,18 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+typedef int int2 __attribute__((ext_vector_type(2)));
+
+struct S {
+ int x;
+};
+
void test_gfx9_fmed3h(global half *out, half a, half b, half c)
{
*out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
}
-void test_mov_dpp(global int* out, int src, int i)
+void test_mov_dpp(global int* out, int src, int i, int2 i2, struct S s, float _Complex fc)
{
*out = __builtin_amdgcn_mov_dpp(src, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp(src, 0, i, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
@@ -22,9 +28,12 @@ void test_mov_dpp(global int* out, int src, int i)
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false, 1); // expected-error{{too many arguments to function call, expected at most 5, have 6}}
*out = __builtin_amdgcn_mov_dpp(out, 0, 0, 0, false); // expected-error{{used type '__global int *__private' where integer or floating point type is required}}
*out = __builtin_amdgcn_mov_dpp("aa", 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_mov_dpp(i2, 0, 0, 0, false); // expected-error{{used type '__private int2' (vector of 2 'int' values) where integer or floating point type is required}}
+ *out = __builtin_amdgcn_mov_dpp(s, 0, 0, 0, false); // expected-error{{used type '__private struct S' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_mov_dpp(fc, 0, 0, 0, false); // expected-error{{used type '__private _Complex float' where integer or floating point type is required}}
}
-void test_update_dpp(global int* out, int arg1, int arg2, int i)
+void test_update_dpp(global int* out, int arg1, int arg2, int i, int2 i2, struct S s, float _Complex fc)
{
*out = __builtin_amdgcn_update_dpp(arg1, arg2, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, i, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
@@ -40,4 +49,10 @@ void test_update_dpp(global int* out, int arg1, int arg2, int i)
*out = __builtin_amdgcn_update_dpp(arg1, out, 0, 0, 0, false); // expected-error{{used type '__global int *__private' where integer or floating point type is required}}
*out = __builtin_amdgcn_update_dpp("aa", arg2, 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
*out = __builtin_amdgcn_update_dpp(arg1, "aa", 0, 0, 0, false); // expected-error{{used type '__constant char[3]' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(i2, arg2, 0, 0, 0, false); // expected-error{{used type '__private int2' (vector of 2 'int' values) where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(arg1, i2, 0, 0, 0, false); // expected-error{{used type '__private int2' (vector of 2 'int' values) where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(s, arg2, 0, 0, 0, false); // expected-error{{used type '__private struct S' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(arg1, s, 0, 0, 0, false); // expected-error{{used type '__private struct S' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(fc, arg2, 0, 0, 0, false); // expected-error{{used type '__private _Complex float' where integer or floating point type is required}}
+ *out = __builtin_amdgcn_update_dpp(arg1, fc, 0, 0, 0, false); // expected-error{{used type '__private _Complex float' where integer or floating point type is required}}
}
More information about the cfe-commits
mailing list