yaxunl created this revision.
Herald added subscribers: tpr, dstuttard, tony-tye, nhaehnle, wdng, kzhuravl.
https://reviews.llvm.org/D30551
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/CodeGen/CGBuiltin.cpp
test/CodeGenOpenCL/builtins-amdgcn.cl
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -235,6 +235,41 @@
*out = __builtin_amdgcn_ds_swizzle(a, 32);
}
+// CHECK-LABEL: @test_ds_permute
+// CHECK: call i32 @llvm.amdgcn.ds.permute.i32(i32 %a, i32 %b)
+void test_ds_permute(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_permute(a, b);
+}
+
+// CHECK-LABEL: @test_ds_bpermute
+// CHECK: call i32 @llvm.amdgcn.ds.bpermute.i32(i32 %a, i32 %b)
+void test_ds_bpermute(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_bpermute(a, b);
+}
+
+// CHECK-LABEL: @test_readfirstlane
+// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a)
+void test_readfirstlane(global int* out, int a)
+{
+ *out = __builtin_amdgcn_readfirstlane(a);
+}
+
+// CHECK-LABEL: @test_readlane
+// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b)
+void test_readlane(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_readlane(a, b);
+}
+
+// CHECK-LABEL: @test_mov_dpp
+// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 %dpp_ctrl, i32
%row_mask, i32 %bank_mask, i1 %bound_ctrl)
+void test_mov_dpp(global int* out, int src, int dpp_ctrl, int row_mask, int
bank_mask, bool bound_ctrl)
+{
+ *out = __builtin_amdgcn_mov_dpp(src, dpp_ctrl, row_mask, bank_mask,
bound_ctrl);
+}
+
// CHECK-LABEL: @test_fcmp_f32
// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
void test_fcmp_f32(global ulong* out, float a, float b)
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -8388,6 +8388,22 @@
case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle);
+ case AMDGPU::BI__builtin_amdgcn_ds_permute:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_permute);
+ case AMDGPU::BI__builtin_amdgcn_ds_bpermute:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_bpermute);
+ case AMDGPU::BI__builtin_amdgcn_readfirstlane:
+ return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_readfirstlane);
+ case AMDGPU::BI__builtin_amdgcn_readlane:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_readlane);
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
+ llvm::SmallVector<llvm::Value *, 5> Args;
+ for (unsigned I = 0; I != 5; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_mov_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
@@ -86,6 +86,11 @@
BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc")
+BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
+BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
+BUILTIN(__builtin_amdgcn_mov_dpp, "iiiiib", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
//===----------------------------------------------------------------------===//
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -235,6 +235,41 @@
*out = __builtin_amdgcn_ds_swizzle(a, 32);
}
+// CHECK-LABEL: @test_ds_permute
+// CHECK: call i32 @llvm.amdgcn.ds.permute.i32(i32 %a, i32 %b)
+void test_ds_permute(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_permute(a, b);
+}
+
+// CHECK-LABEL: @test_ds_bpermute
+// CHECK: call i32 @llvm.amdgcn.ds.bpermute.i32(i32 %a, i32 %b)
+void test_ds_bpermute(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_bpermute(a, b);
+}
+
+// CHECK-LABEL: @test_readfirstlane
+// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a)
+void test_readfirstlane(global int* out, int a)
+{
+ *out = __builtin_amdgcn_readfirstlane(a);
+}
+
+// CHECK-LABEL: @test_readlane
+// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b)
+void test_readlane(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_readlane(a, b);
+}
+
+// CHECK-LABEL: @test_mov_dpp
+// CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 %dpp_ctrl, i32 %row_mask, i32 %bank_mask, i1 %bound_ctrl)
+void test_mov_dpp(global int* out, int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl)
+{
+ *out = __builtin_amdgcn_mov_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+}
+
// CHECK-LABEL: @test_fcmp_f32
// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
void test_fcmp_f32(global ulong* out, float a, float b)
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -8388,6 +8388,22 @@
case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle);
+ case AMDGPU::BI__builtin_amdgcn_ds_permute:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_permute);
+ case AMDGPU::BI__builtin_amdgcn_ds_bpermute:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_bpermute);
+ case AMDGPU::BI__builtin_amdgcn_readfirstlane:
+ return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_readfirstlane);
+ case AMDGPU::BI__builtin_amdgcn_readlane:
+ return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_readlane);
+ case AMDGPU::BI__builtin_amdgcn_mov_dpp: {
+ llvm::SmallVector<llvm::Value *, 5> Args;
+ for (unsigned I = 0; I != 5; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_mov_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
@@ -86,6 +86,11 @@
BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc")
+BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
+BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
+BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
+BUILTIN(__builtin_amdgcn_mov_dpp, "iiiiib", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
