r332848 - [AMDGPU] fixes for lds f32 builtins
Daniil Fukalov via cfe-commits
cfe-commits at lists.llvm.org
Mon May 21 09:18:07 PDT 2018
Author: dfukalov
Date: Mon May 21 09:18:07 2018
New Revision: 332848
URL: http://llvm.org/viewvc/llvm-project?rev=332848&view=rev
Log:
[AMDGPU] fixes for lds f32 builtins
1. added restrictions to memory scope, order and volatile parameters
2. added custom processing for these builtins - currently is not used code,
needed to switch off GCCBuiltin link to the builtins (ongoing change to llvm
tree)
3. builtins renamed as requested
Differential Revision: https://reviews.llvm.org/D43281
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=332848&r1=332847&r2=332848&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon May 21 09:18:07 2018
@@ -93,9 +93,9 @@ BUILTIN(__builtin_amdgcn_ds_bpermute, "i
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_faddf, "ff*fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fminf, "ff*fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*fIiIiIb", "n")
//===----------------------------------------------------------------------===//
// VI+ only builtins.
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=332848&r1=332847&r2=332848&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon May 21 09:18:07 2018
@@ -10088,6 +10088,49 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
CI->setConvergent();
return CI;
}
+ case AMDGPU::BI__builtin_amdgcn_ds_faddf:
+ case AMDGPU::BI__builtin_amdgcn_ds_fminf:
+ case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: {
+ llvm::SmallVector<llvm::Value *, 5> Args;
+ for (unsigned I = 0; I != 5; ++I)
+ Args.push_back(EmitScalarExpr(E->getArg(I)));
+ const llvm::Type *PtrTy = Args[0]->getType();
+ // check pointer parameter
+ if (!PtrTy->isPointerTy() ||
+ E->getArg(0)
+ ->getType()
+ ->getPointeeType()
+ .getQualifiers()
+ .getAddressSpace() != LangAS::opencl_local ||
+ !PtrTy->getPointerElementType()->isFloatTy()) {
+ CGM.Error(E->getArg(0)->getLocStart(),
+ "parameter should have type \"local float*\"");
+ return nullptr;
+ }
+ // check float parameter
+ if (!Args[1]->getType()->isFloatTy()) {
+ CGM.Error(E->getArg(1)->getLocStart(),
+ "parameter should have type \"float\"");
+ return nullptr;
+ }
+
+ Intrinsic::ID ID;
+ switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_ds_faddf:
+ ID = Intrinsic::amdgcn_ds_fadd;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_ds_fminf:
+ ID = Intrinsic::amdgcn_ds_fmin;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
+ 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:
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=332848&r1=332847&r2=332848&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Mon May 21 09:18:07 2018
@@ -91,18 +91,18 @@ void test_mov_dpp(global int* out, int s
// 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) {
- *out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, false);
+void test_ds_faddf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_faddf(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) {
- *out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, false);
+void test_ds_fminf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_fminf(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) {
- *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
+void test_ds_fmaxf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
}
Modified: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl?rev=332848&r1=332847&r2=332848&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (original)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Mon May 21 09:18:07 2018
@@ -102,3 +102,20 @@ void test_mov_dpp2(global int* out, int
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
}
+void test_ds_faddf(local float *out, float src, int a) {
+ *out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+}
+
+void test_ds_fminf(local float *out, float src, int a) {
+ *out = __builtin_amdgcn_ds_fminf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_fminf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+}
+
+void test_ds_fmaxf(local float *out, float src, int a) {
+ *out = __builtin_amdgcn_ds_fmaxf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+ *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+}
More information about the cfe-commits
mailing list