From 03bdd8f797761dd10927f283f51e612db9e9f788 Mon Sep 17 00:00:00 2001 From: Changpeng Fang Date: Thu, 18 Aug 2016 22:04:54 +0000 Subject: [PATCH] AMDGPU: Add clang builtin for ds_swizzle. Summary: int __builtin_amdgcn_ds_swizzle (int a, int imm); while imm is a constant. Differential Revision: http://reviews.llvm.org/D23682 llvm-svn: 279165 --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 + clang/lib/CodeGen/CGBuiltin.cpp | 3 +++ clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl | 4 ++++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 7 +++++++ 4 files changed, 15 insertions(+) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index b4314e6..0e54128 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -76,6 +76,7 @@ BUILTIN(__builtin_amdgcn_sicmp, "LUiiiIi", "nc") BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiLiIi", "nc") BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc") BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc") +BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc") //===----------------------------------------------------------------------===// // VI+ only builtins. diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 87a825d..c06fcf7 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -7652,6 +7652,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::Value *Src3ToBool = Builder.CreateIsNotNull(Src3); return Builder.CreateCall(F, {Src0, Src1, Src2, Src3ToBool}); } + + case AMDGPU::BI__builtin_amdgcn_ds_swizzle: + return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle); case AMDGPU::BI__builtin_amdgcn_div_fixup: case AMDGPU::BI__builtin_amdgcn_div_fixupf: return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl index 5c67666..5513c5a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl @@ -48,3 +48,7 @@ void test_fcmp_f64(global ulong* out, double a, double b, uint c) *out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}} } +void test_ds_swizzle(global int* out, int a, int b) +{ + *out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}} +} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 2347bc8..3280f19 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -228,6 +228,13 @@ void test_uicmp_i64(global ulong* out, ulong a, ulong b) *out = __builtin_amdgcn_uicmpl(a, b, 30+5); } +// CHECK-LABEL: @test_ds_swizzle +// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32) +void test_ds_swizzle(global int* out, int a) +{ + *out = __builtin_amdgcn_ds_swizzle(a, 32); +} + // 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) -- 2.7.4