[PATCH] D18561: [NVPTX] Set NVPTXTTI::getInliningThresholdMultiplier to 5.

[Justin Lebar via llvm-commits]

jlebar created this revision.
jlebar added a reviewer: chandlerc.
jlebar added subscribers: tra, llvm-commits.
Herald added a subscriber: jholewinski.

Calls on NVPTX are unusually expensive (for one thing, lots of state
needs to be saved to memory, which is slow), so make the inlininer much
more aggressive.

http://reviews.llvm.org/D18561

Files:
  lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
  lib/Target/NVPTX/NVPTXTargetTransformInfo.h

Index: lib/Target/NVPTX/NVPTXTargetTransformInfo.h
===================================================================
--- lib/Target/NVPTX/NVPTXTargetTransformInfo.h
+++ lib/Target/NVPTX/NVPTXTargetTransformInfo.h
@@ -52,6 +52,8 @@
 
   bool isSourceOfDivergence(const Value *V);
 
+  double getInliningThresholdMultiplier(const Function *Caller);
+
   int getArithmeticInstrCost(
       unsigned Opcode, Type *Ty,
       TTI::OperandValueKind Opd1Info = TTI::OK_AnyValue,
Index: lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
+++ lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp
@@ -46,6 +46,13 @@
   }
 }
 
+double
+NVPTXTTIImpl::getInliningThresholdMultiplier(const Function * /* Caller */) {
+  // Increase the inlining cost threshold by a factor of 5, reflecting that
+  // calls are particularly expensive in NVPTX.
+  return 5;
+}
+
 bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) {
   // Without inter-procedural analysis, we conservatively assume that arguments
   // to __device__ functions are divergent.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D18561.51938.patch
Type: text/x-patch
Size: 1140 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160329/4519525b/attachment.bin>