[PATCH] D43281: [AMDGPU] fixes for lds f32 builtins
Daniil Fukalov via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Wed Feb 14 03:16:31 PST 2018
dfukalov created this revision.
dfukalov added reviewers: b-sumner, arsenm.
dfukalov added a project: AMDGPU.
Herald added subscribers: cfe-commits, t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng, kzhuravl.
1. removed addrspace 3 specifications from builtins description strings since it's not target addrspaces
2. added custom processing for these builtins
Repository:
rC Clang
https://reviews.llvm.org/D43281
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/CodeGen/CGBuiltin.cpp
test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Index: test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -91,18 +91,18 @@
// CHECK-LABEL: @test_ds_fadd
// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fadd(__attribute__((address_space(3))) float *out, float src) {
+void test_ds_fadd(local float *out, float src) {
*out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, false);
}
// CHECK-LABEL: @test_ds_fmin
// CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmin(__attribute__((address_space(3))) float *out, float src) {
+void test_ds_fmin(local float *out, float src) {
*out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, false);
}
// CHECK-LABEL: @test_ds_fmax
// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmax(__attribute__((address_space(3))) float *out, float src) {
+void test_ds_fmax(local float *out, float src) {
*out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -9860,6 +9860,29 @@
CI->setConvergent();
return CI;
}
+ case AMDGPU::BI__builtin_amdgcn_ds_fadd:
+ case AMDGPU::BI__builtin_amdgcn_ds_fmin:
+ case AMDGPU::BI__builtin_amdgcn_ds_fmax: {
+ llvm::SmallVector<llvm::Value *, 5> Args;
+ for (unsigned I = 0; I != 5; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ Intrinsic::ID ID;
+ switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_ds_fadd:
+ ID = Intrinsic::amdgcn_ds_fadd;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_ds_fmin:
+ ID = Intrinsic::amdgcn_ds_fmin;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_ds_fmax:
+ ID = Intrinsic::amdgcn_ds_fmax;
+ break;
+ default:
+ llvm_unreachable("Unknown BuiltinID");
+ }
+ Value *F = CGM.getIntrinsic(ID);
+ return Builder.CreateCall(F, Args);
+ }
// amdgcn workitem
case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -93,9 +93,9 @@
BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
-BUILTIN(__builtin_amdgcn_ds_fadd, "ff*3fiib", "n")
-BUILTIN(__builtin_amdgcn_ds_fmin, "ff*3fiib", "n")
-BUILTIN(__builtin_amdgcn_ds_fmax, "ff*3fiib", "n")
+BUILTIN(__builtin_amdgcn_ds_fadd, "ff*fiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmin, "ff*fiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmax, "ff*fiIiIb", "n")
//===----------------------------------------------------------------------===//
// VI+ only builtins.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D43281.134185.patch
Type: text/x-patch
Size: 3075 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20180214/df2a0cd9/attachment.bin>
More information about the llvm-commits
mailing list