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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@279165 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index b4314e6b0fff0b157b43d02c909478d500606878..0e54128080f6308255d01276b59adb75bf3a0a1a 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/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/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp
index 87a825d46ade7f535b769cab6585e60173fb3332..c06fcf73f379f5360ae4d84a79e9a8d81f1675b7 100644 (file)
--- a/lib/CodeGen/CGBuiltin.cpp
+++ b/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/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl
index 5c676664c1a14071efdb7ab46d4393de93ef2f90..5513c5a9587aaf8aa5fb9a13e4884dec747bdd6b 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn-error.cl
+++ b/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/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index 2347bc8e41c81fb0ef73304772025bca693694a3..3280f1907617ce425235f08405d1afe7206b56e2 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/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)