yaxunl created this revision.
yaxunl added reviewers: kzhuravl, b-sumner, arsenm.
Herald added subscribers: t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely.
https://reviews.llvm.org/D52320
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/CodeGen/CGBuiltin.cpp
test/CodeGenOpenCL/builtins-amdgcn-vi.cl
test/SemaOpenCL/builtins-amdgcn-error.cl
Index: test/SemaOpenCL/builtins-amdgcn-error.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error.cl
+++ test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -102,6 +102,15 @@
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument
to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
}
+void test_update_dpp2(global int* out, int a, int b, int c, int d, int e, bool
f)
+{
+ *out = __builtin_amdgcn_update_dpp(a, b, 0, 0, 0, false);
+ *out = __builtin_amdgcn_update_dpp(a, 0, c, 0, 0, false); // expected-error
{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, d, 0, false); // expected-error
{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, e, false); // expected-error
{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, 0, f); // expected-error
{{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+}
+
void test_ds_faddf(local float *out, float src, int a) {
*out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error
{{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error
{{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
Index: test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -96,6 +96,13 @@
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
}
+// CHECK-LABEL: @test_update_dpp
+// 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)
+{
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
+}
+
// CHECK-LABEL: @test_ds_fadd
// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float
%src, i32 0, i32 0, i1 false)
void test_ds_faddf(local float *out, float src) {
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -11310,6 +11310,14 @@
Args[0]->getType());
return Builder.CreateCall(F, Args);
}
+ case AMDGPU::BI__builtin_amdgcn_update_dpp: {
+ llvm::SmallVector<llvm::Value *, 6> Args;
+ for (unsigned I = 0; I != 6; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Value *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
+ return Builder.CreateCall(F, Args);
+ }
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
case AMDGPU::BI__builtin_amdgcn_div_fixuph:
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -122,6 +122,7 @@
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "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_s_dcache_wb, "v", "n", "vi-insts")
//===----------------------------------------------------------------------===//
Index: test/SemaOpenCL/builtins-amdgcn-error.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error.cl
+++ test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -102,6 +102,15 @@
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
}
+void test_update_dpp2(global int* out, int a, int b, int c, int d, int e, bool f)
+{
+ *out = __builtin_amdgcn_update_dpp(a, b, 0, 0, 0, false);
+ *out = __builtin_amdgcn_update_dpp(a, 0, c, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, d, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, e, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+ *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, 0, f); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+}
+
void test_ds_faddf(local float *out, float src, int a) {
*out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
Index: test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -96,6 +96,13 @@
*out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
}
+// CHECK-LABEL: @test_update_dpp
+// 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)
+{
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
+}
+
// CHECK-LABEL: @test_ds_fadd
// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
void test_ds_faddf(local float *out, float src) {
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -11310,6 +11310,14 @@
Args[0]->getType());
return Builder.CreateCall(F, Args);
}
+ case AMDGPU::BI__builtin_amdgcn_update_dpp: {
+ llvm::SmallVector<llvm::Value *, 6> Args;
+ for (unsigned I = 0; I != 6; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Value *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_update_dpp, Args[0]->getType());
+ return Builder.CreateCall(F, Args);
+ }
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
case AMDGPU::BI__builtin_amdgcn_div_fixuph:
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -122,6 +122,7 @@
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "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_s_dcache_wb, "v", "n", "vi-insts")
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
