[PATCH] D23682: AMDGPU: Add Clang builtin intrinsics for ds_swizzle

[Phabricator via llvm-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL279165: AMDGPU: Add clang builtin for ds_swizzle. (authored by chfang).

Changed prior to commit:
  https://reviews.llvm.org/D23682?vs=68604&id=68624#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D23682

Files:
  cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
  cfe/trunk/lib/CodeGen/CGBuiltin.cpp
  cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
  cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

Index: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp
@@ -7652,6 +7652,9 @@
     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);
Index: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
@@ -76,6 +76,7 @@
 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.
Index: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -228,6 +228,13 @@
   *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)
Index: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
@@ -48,3 +48,7 @@
   *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}}
+}


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D23682.68624.patch
Type: text/x-patch
Size: 2432 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160818/0830c91e/attachment.bin>