Author: arsenm
Date: Mon Jan 28 15:59:18 2019
New Revision: 352443
URL: http://llvm.org/viewvc/llvm-project?rev=352443&view=rev
Log:
AMDGPU: Add ds append/consume builtins
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=352443&r1=352442&r2=352443&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Jan 28 15:59:18 2019
@@ -98,6 +98,8 @@ BUILTIN(__builtin_amdgcn_fmed3f, "ffff",
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.
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=352443&r1=352442&r2=352443&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon Jan 28 15:59:18 2019
@@ -12478,6 +12478,14 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
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"));
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=352443&r1=352442&r2=352443&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Mon Jan 28 15:59:18 2019
@@ -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 }
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
