]> granicus.if.org Git - clang/commitdiff
AMDGPU: Add ds append/consume builtins
authorMatt Arsenault <Matthew.Arsenault@amd.com>
Mon, 28 Jan 2019 23:59:18 +0000 (23:59 +0000)
committerMatt Arsenault <Matthew.Arsenault@amd.com>
Mon, 28 Jan 2019 23:59:18 +0000 (23:59 +0000)
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@352443 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/Basic/BuiltinsAMDGPU.def
lib/CodeGen/CGBuiltin.cpp
test/CodeGenOpenCL/builtins-amdgcn.cl

index 547870447dd52805482f2fc17db428f3d43d02c4..fa3db618be7cedbbbf6b20a80868366d8fcdcc7a 100644 (file)
@@ -98,6 +98,8 @@ BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_append, "ii*3", "n")
+BUILTIN(__builtin_amdgcn_ds_consume, "ii*3", "n")
 
 //===----------------------------------------------------------------------===//
 // CI+ only builtins.
index 195b0c078bd22a4881b9b56bc8374156fbe33a8a..fa4d5126433310716a15a3873e16c7d94cb20867 100644 (file)
@@ -12478,6 +12478,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_fmed3f:
   case AMDGPU::BI__builtin_amdgcn_fmed3h:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fmed3);
+  case AMDGPU::BI__builtin_amdgcn_ds_append:
+  case AMDGPU::BI__builtin_amdgcn_ds_consume: {
+    Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?
+      Intrinsic::amdgcn_ds_append : Intrinsic::amdgcn_ds_consume;
+    Value *Src0 = EmitScalarExpr(E->getArg(0));
+    Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() });
+    return Builder.CreateCall(F, { Src0, Builder.getFalse() });
+  }
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));
index dc7f480209af33fa7eea3f9dfd798adda739665e..6b7ea52dab0bc442ef65bcd0cbf6007d0742bba6 100644 (file)
@@ -536,6 +536,18 @@ void test_s_getpc(global ulong* out)
   *out = __builtin_amdgcn_s_getpc();
 }
 
+// CHECK-LABEL: @test_ds_append_lds(
+// CHECK: call i32 @llvm.amdgcn.ds.append.p3i32(i32 addrspace(3)* %ptr, i1 false)
+kernel void test_ds_append_lds(global int* out, local int* ptr) {
+  *out = __builtin_amdgcn_ds_append(ptr);
+}
+
+// CHECK-LABEL: @test_ds_consume_lds(
+// CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false)
+kernel void test_ds_consume_lds(global int* out, local int* ptr) {
+  *out = __builtin_amdgcn_ds_consume(ptr);
+}
+
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
 // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }