From 583567bc90a3c3df6094f2ad5d64de451fc645c5 Mon Sep 17 00:00:00 2001
From: Eric Niebler <eniebler@nvidia.com>
Date: Fri, 4 Oct 2024 12:49:39 -0700
Subject: [PATCH] avoid gcc optimizer bug by not force inlining part of
 `thrust::transform` (#2509)

---
 thrust/thrust/system/cuda/detail/transform.h | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/thrust/thrust/system/cuda/detail/transform.h b/thrust/thrust/system/cuda/detail/transform.h
index 9e1d0b2a318..1926fb62473 100644
--- a/thrust/thrust/system/cuda/detail/transform.h
+++ b/thrust/thrust/system/cuda/detail/transform.h
@@ -178,8 +178,10 @@ struct binary_transform_f<InputIt1, InputIt2, OutputIt, no_stencil_tag, Transfor
   }
 }; // struct binary_transform_f
 
+// EAN 2024-10-04: when force-inlined, gcc's optimizer will generate bad code
+// for this function:
 template <class Policy, class InputIt, class Size, class OutputIt, class StencilIt, class TransformOp, class Predicate>
-OutputIt THRUST_FUNCTION unary(
+OutputIt _CCCL_HOST_DEVICE inline unary(
   Policy& policy,
   InputIt items,
   OutputIt result,
@@ -200,6 +202,8 @@ OutputIt THRUST_FUNCTION unary(
   return result + num_items;
 }
 
+// EAN 2024-10-04: when force-inlined, gcc's optimizer will generate bad code
+// for this function:
 template <class Policy,
           class InputIt1,
           class InputIt2,
@@ -208,7 +212,7 @@ template <class Policy,
           class StencilIt,
           class TransformOp,
           class Predicate>
-OutputIt THRUST_FUNCTION binary(
+OutputIt _CCCL_HOST_DEVICE inline binary(
   Policy& policy,
   InputIt1 items1,
   InputIt2 items2,