Skip to content

Commit af4fc56

Browse files
rampitecronlieb
rampitec
authored and
ronlieb
committed
[AMDGPU] Allow overload of __builtin_amdgcn_mov/update_dpp (llvm#112447)
We need to support 64-bit data types (intrinsics do support it). We are also silently converting FP to integer argument now, also fixed. Change-Id: I2a78e8e687884625f88e06da117fe7b8cb20dd01
1 parent e51df59 commit af4fc56

File tree

5 files changed

+220
-13
lines changed

5 files changed

+220
-13
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -224,8 +224,8 @@ TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
224224
TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
225225
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
226226
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "WUi", "n", "s-memrealtime")
227-
TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
228-
TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp")
227+
TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nct", "dpp")
228+
TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nct", "dpp")
229229
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "gfx8-insts")
230230
TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
231231

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 24 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -19041,15 +19041,32 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1904119041
ASTContext::GetBuiltinTypeError Error;
1904219042
getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);
1904319043
assert(Error == ASTContext::GE_None && "Should not codegen an error");
19044+
llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());
19045+
unsigned Size = DataTy->getPrimitiveSizeInBits();
19046+
llvm::Type *IntTy =
19047+
llvm::IntegerType::get(Builder.getContext(), std::max(Size, 32u));
19048+
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, IntTy);
19049+
assert(E->getNumArgs() == 5 || E->getNumArgs() == 6);
19050+
bool InsertOld = E->getNumArgs() == 5;
19051+
if (InsertOld)
19052+
Args.push_back(llvm::PoisonValue::get(IntTy));
1904419053
for (unsigned I = 0; I != E->getNumArgs(); ++I) {
19045-
Args.push_back(EmitScalarOrConstFoldImmArg(ICEArguments, I, E));
19054+
llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);
19055+
if (I <= !InsertOld && Size < 32) {
19056+
if (!DataTy->isIntegerTy())
19057+
V = Builder.CreateBitCast(
19058+
V, llvm::IntegerType::get(Builder.getContext(), Size));
19059+
V = Builder.CreateZExtOrBitCast(V, IntTy);
19060+
}
19061+
llvm::Type *ExpTy =
19062+
F->getFunctionType()->getFunctionParamType(I + InsertOld);
19063+
Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy));
1904619064
}
19047-
assert(Args.size() == 5 || Args.size() == 6);
19048-
if (Args.size() == 5)
19049-
Args.insert(Args.begin(), llvm::PoisonValue::get(Args[0]->getType()));
19050-
Function *F =
19051-
CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
19052-
return Builder.CreateCall(F, Args);
19065+
Value *V = Builder.CreateCall(F, Args);
19066+
if (Size < 32 && !DataTy->isIntegerTy())
19067+
V = Builder.CreateTrunc(
19068+
V, llvm::IntegerType::get(Builder.getContext(), Size));
19069+
return Builder.CreateTruncOrBitCast(V, DataTy);
1905319070
}
1905419071
case AMDGPU::BI__builtin_amdgcn_permlane16:
1905519072
case AMDGPU::BI__builtin_amdgcn_permlanex16:

‎clang/lib/Sema/SemaAMDGPU.cpp

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -63,6 +63,44 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
6363
OrderIndex = 0;
6464
ScopeIndex = 1;
6565
break;
66+
case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
67+
if (SemaRef.checkArgCountRange(TheCall, 5, 5))
68+
return true;
69+
Expr *ValArg = TheCall->getArg(0);
70+
QualType Ty = ValArg->getType();
71+
// TODO: Vectors can also be supported.
72+
if (!Ty->isArithmeticType() || Ty->isAnyComplexType()) {
73+
SemaRef.Diag(ValArg->getBeginLoc(),
74+
diag::err_typecheck_cond_expect_int_float)
75+
<< Ty << ValArg->getSourceRange();
76+
return true;
77+
}
78+
return false;
79+
}
80+
case AMDGPU::BI__builtin_amdgcn_update_dpp: {
81+
if (SemaRef.checkArgCountRange(TheCall, 6, 6))
82+
return true;
83+
Expr *Args[2];
84+
QualType ArgTys[2];
85+
for (unsigned I = 0; I != 2; ++I) {
86+
Args[I] = TheCall->getArg(I);
87+
ArgTys[I] = Args[I]->getType();
88+
// TODO: Vectors can also be supported.
89+
if (!ArgTys[I]->isArithmeticType() || ArgTys[I]->isAnyComplexType()) {
90+
SemaRef.Diag(Args[I]->getBeginLoc(),
91+
diag::err_typecheck_cond_expect_int_float)
92+
<< ArgTys[I] << Args[I]->getSourceRange();
93+
return true;
94+
}
95+
}
96+
if (ArgTys[0] != ArgTys[1]) {
97+
SemaRef.Diag(Args[1]->getBeginLoc(),
98+
diag::err_typecheck_call_different_arg_types)
99+
<< ArgTys[0] << ArgTys[1];
100+
return true;
101+
}
102+
return false;
103+
}
66104
default:
67105
return false;
68106
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

Lines changed: 106 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -102,20 +102,122 @@ void test_s_dcache_wb()
102102
__builtin_amdgcn_s_dcache_wb();
103103
}
104104

105-
// CHECK-LABEL: @test_mov_dpp
105+
// CHECK-LABEL: @test_mov_dpp_int
106106
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %src, i32 0, i32 0, i32 0, i1 false)
107-
void test_mov_dpp(global int* out, int src)
107+
void test_mov_dpp_int(global int* out, int src)
108108
{
109109
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
110110
}
111111

112-
// CHECK-LABEL: @test_update_dpp
112+
// CHECK-LABEL: @test_mov_dpp_long
113+
// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %x, i32 257, i32 15, i32 15, i1 false)
114+
// CHECK-NEXT: store i64 %0,
115+
void test_mov_dpp_long(long x, global long *p) {
116+
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
117+
}
118+
119+
// CHECK-LABEL: @test_mov_dpp_float
120+
// CHECK: %0 = bitcast float %x to i32
121+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
122+
// CHECK-NEXT: store i32 %1,
123+
void test_mov_dpp_float(float x, global float *p) {
124+
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
125+
}
126+
127+
// CHECK-LABEL: @test_mov_dpp_double
128+
// CHECK: %0 = bitcast double %x to i64
129+
// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 poison, i64 %0, i32 257, i32 15, i32 15, i1 false)
130+
// CHECK-NEXT: store i64 %1,
131+
void test_mov_dpp_double(double x, global double *p) {
132+
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
133+
}
134+
135+
// CHECK-LABEL: @test_mov_dpp_short
136+
// CHECK: %0 = zext i16 %x to i32
137+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
138+
// CHECK-NEXT: %2 = trunc i32 %1 to i16
139+
// CHECK-NEXT: store i16 %2,
140+
void test_mov_dpp_short(short x, global short *p) {
141+
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
142+
}
143+
144+
// CHECK-LABEL: @test_mov_dpp_char
145+
// CHECK: %0 = zext i8 %x to i32
146+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %0, i32 257, i32 15, i32 15, i1 false)
147+
// CHECK-NEXT: %2 = trunc i32 %1 to i8
148+
// CHECK-NEXT: store i8 %2,
149+
void test_mov_dpp_char(char x, global char *p) {
150+
*p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
151+
}
152+
153+
// CHECK-LABEL: @test_mov_dpp_half
154+
// CHECK: %0 = load i16,
155+
// CHECK: %1 = zext i16 %0 to i32
156+
// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %1, i32 257, i32 15, i32 15, i1 false)
157+
// CHECK-NEXT: %3 = trunc i32 %2 to i16
158+
// CHECK-NEXT: store i16 %3,
159+
void test_mov_dpp_half(half *x, global half *p) {
160+
*p = __builtin_amdgcn_mov_dpp(*x, 0x101, 0xf, 0xf, 0);
161+
}
162+
163+
// CHECK-LABEL: @test_update_dpp_int
113164
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
114-
void test_update_dpp(global int* out, int arg1, int arg2)
165+
void test_update_dpp_int(global int* out, int arg1, int arg2)
115166
{
116167
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
117168
}
118169

170+
// CHECK-LABEL: @test_update_dpp_long
171+
// CHECK: %0 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %x, i64 %x, i32 257, i32 15, i32 15, i1 false)
172+
// CHECk-NEXT: store i64 %0,
173+
void test_update_dpp_long(long x, global long *p) {
174+
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
175+
}
176+
177+
// CHECK-LABEL: @test_update_dpp_float
178+
// CHECK: %0 = bitcast float %x to i32
179+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
180+
// CHECK-NEXT: store i32 %1,
181+
void test_update_dpp_float(float x, global float *p) {
182+
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
183+
}
184+
185+
// CHECK-LABEL: @test_update_dpp_double
186+
// CHECK: %0 = bitcast double %x to i64
187+
// CHECK-NEXT: %1 = tail call{{.*}} i64 @llvm.amdgcn.update.dpp.i64(i64 %0, i64 %0, i32 257, i32 15, i32 15, i1 false)
188+
// CHECK-NEXT: store i64 %1,
189+
void test_update_dpp_double(double x, global double *p) {
190+
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
191+
}
192+
193+
// CHECK-LABEL: @test_update_dpp_short
194+
// CHECK: %0 = zext i16 %x to i32
195+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
196+
// CHECK-NEXT: %2 = trunc i32 %1 to i16
197+
// CHECK-NEXT: store i16 %2,
198+
void test_update_dpp_short(short x, global short *p) {
199+
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
200+
}
201+
202+
// CHECK-LABEL: @test_update_dpp_char
203+
// CHECK: %0 = zext i8 %x to i32
204+
// CHECK-NEXT: %1 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %0, i32 %0, i32 257, i32 15, i32 15, i1 false)
205+
// CHECK-NEXT: %2 = trunc i32 %1 to i8
206+
// CHECK-NEXT: store i8 %2,
207+
void test_update_dpp_char(char x, global char *p) {
208+
*p = __builtin_amdgcn_update_dpp(x, x, 0x101, 0xf, 0xf, 0);
209+
}
210+
211+
// CHECK-LABEL: @test_update_dpp_half
212+
// CHECK: %0 = load i16,
213+
// CHECK: %1 = zext i16 %0 to i32
214+
// CHECK-NEXT: %2 = tail call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %1, i32 %1, i32 257, i32 15, i32 15, i1 false)
215+
// CHECK-NEXT: %3 = trunc i32 %2 to i16
216+
// CHECK-NEXT: store i16 %3,
217+
void test_update_dpp_half(half *x, global half *p) {
218+
*p = __builtin_amdgcn_update_dpp(*x, *x, 0x101, 0xf, 0xf, 0);
219+
}
220+
119221
// CHECK-LABEL: @test_ds_fadd
120222
// CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}
121223
// CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}}

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,57 @@
33

44
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
55

6+
typedef int int2 __attribute__((ext_vector_type(2)));
7+
8+
struct S {
9+
int x;
10+
};
11+
612
void test_gfx9_fmed3h(global half *out, half a, half b, half c)
713
{
814
*out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
915
}
16+
17+
void test_mov_dpp(global int* out, int src, int i, int2 i2, struct S s, float _Complex fc)
18+
{
19+
*out = __builtin_amdgcn_mov_dpp(src, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
20+
*out = __builtin_amdgcn_mov_dpp(src, 0, i, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
21+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, i, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
22+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, i); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
23+
*out = __builtin_amdgcn_mov_dpp(src, 0.1, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
24+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0.1, 0, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
25+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0.1, false); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
26+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, 0.1); // expected-error{{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
27+
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0); // expected-error{{too few arguments to function call, expected 5, have 4}}
28+
*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}}
29+
*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}}
30+
*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}}
31+
*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}}
32+
*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}}
33+
*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}}
34+
}
35+
36+
void test_update_dpp(global int* out, int arg1, int arg2, int i, int2 i2, long l, struct S s, float _Complex fc)
37+
{
38+
*out = __builtin_amdgcn_update_dpp(arg1, arg2, i, 0, 0, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
39+
*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+
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, i, false); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
41+
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, i); // expected-error{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
42+
*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}}
43+
*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}}
44+
*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}}
45+
*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}}
46+
*out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0); // expected-error{{too few arguments to function call, expected 6, have 5}}
47+
*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}}
48+
*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}}
49+
*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}}
50+
*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}}
51+
*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}}
52+
*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}}
53+
*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}}
54+
*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}}
55+
*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}}
56+
*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}}
57+
*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}}
58+
*out = __builtin_amdgcn_update_dpp(i, l, 0, 0, 0, false); // expected-error{{arguments are of different types ('__private int' vs '__private long')}}
59+
}

0 commit comments

Comments
 (0)