[PATCH] D43281: [AMDGPU] fixes for lds f32 builtins

[Daniil Fukalov via Phabricator via llvm-commits]

dfukalov updated this revision to Diff 136752.
dfukalov edited the summary of this revision.
dfukalov set the repository for this revision to rC Clang.
dfukalov added a comment.

addrspace specifications are kept in descriptions strings


Repository:
  rC Clang

https://reviews.llvm.org/D43281

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


Index: test/SemaOpenCL/builtins-amdgcn-error.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error.cl
+++ test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -102,3 +102,20 @@
   *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_fadd(__attribute__((address_space(3))) float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_fadd(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fadd' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fadd(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fadd' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fadd' must be a constant integer}}
+}
+
+void test_ds_fmin(__attribute__((address_space(3))) float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_fmin(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmin' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmin(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmin' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmin' must be a constant integer}}
+}
+
+void test_ds_fmax(__attribute__((address_space(3))) float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_fmax(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmax' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmax(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmax' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmax' must be a constant integer}}
+}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -9860,6 +9860,49 @@
     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)));
+    const llvm::Type *PtrTy = Args[0]->getType();
+    // check pointer parameter
+    if (!PtrTy->isPointerTy() ||
+        LangAS::opencl_local != E->getArg(0)
+                                    ->getType()
+                                    ->getPointeeType()
+                                    .getQualifiers()
+                                    .getAddressSpace() ||
+        !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_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*3fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmin, "ff*3fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmax, "ff*3fIiIiIb", "n")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D43281.136752.patch
Type: text/x-patch
Size: 4611 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20180302/d04f7d98/attachment.bin>