r264008 - [sema] [CUDA] Use std algorithms in EraseUnwantedCUDAMatchesImpl.

Justin Lebar via cfe-commits cfe-commits at lists.llvm.org
Thu Jun 30 15:27:37 PDT 2016


I think this code is accidentally OK, because the list of overloads
that callers pass in in should have at most two elements.  A given
overload set may have both a __host__ and a __device__ function, or
alternatively it may have a __host__ __device__ function, but you
can't have all three.

There is, however, a substantial gap in our testing, which I
accidentally created when I refactored a test, which causes us never
to call this function with even two overload candidates during tests.
So I'm fixing that, and I'll also fix the function to call the correct
erase() so that it does the right thing in general.

Thanks again.

On Thu, Jun 30, 2016 at 1:43 PM, Justin Lebar <jlebar at google.com> wrote:
> Interestingly all the clang tests pass with that whole line commented
> out.  So something *really* seems missing here.
>
> Thank you for finding this.
>
> On Thu, Jun 30, 2016 at 5:08 AM, Benjamin Kramer <benny.kra at gmail.com> wrote:
>> On Tue, Mar 22, 2016 at 1:09 AM, Justin Lebar via cfe-commits
>> <cfe-commits at lists.llvm.org> wrote:
>>> Author: jlebar
>>> Date: Mon Mar 21 19:09:25 2016
>>> New Revision: 264008
>>>
>>> URL: http://llvm.org/viewvc/llvm-project?rev=264008&view=rev
>>> Log:
>>> [sema] [CUDA] Use std algorithms in EraseUnwantedCUDAMatchesImpl.
>>>
>>> Summary: NFC
>>>
>>> Reviewers: tra
>>>
>>> Subscribers: cfe-commits
>>>
>>> Differential Revision: http://reviews.llvm.org/D18327
>>>
>>> Modified:
>>>     cfe/trunk/lib/Sema/SemaCUDA.cpp
>>>
>>> Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
>>> URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=264008&r1=264007&r2=264008&view=diff
>>> ==============================================================================
>>> --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
>>> +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Mon Mar 21 19:09:25 2016
>>> @@ -210,31 +210,28 @@ bool Sema::CheckCUDATarget(const Functio
>>>    return false;
>>>  }
>>>
>>> -template <typename T, typename FetchDeclFn>
>>> -static void EraseUnwantedCUDAMatchesImpl(Sema &S, const FunctionDecl *Caller,
>>> -                                         llvm::SmallVectorImpl<T> &Matches,
>>> -                                         FetchDeclFn FetchDecl) {
>>> +template <typename T>
>>> +static void EraseUnwantedCUDAMatchesImpl(
>>> +    Sema &S, const FunctionDecl *Caller, llvm::SmallVectorImpl<T> &Matches,
>>> +    std::function<const FunctionDecl *(const T &)> FetchDecl) {
>>>    assert(S.getLangOpts().CUDATargetOverloads &&
>>>           "Should not be called w/o enabled target overloads.");
>>>    if (Matches.size() <= 1)
>>>      return;
>>>
>>> +  // Gets the CUDA function preference for a call from Caller to Match.
>>> +  auto GetCFP = [&](const T &Match) {
>>> +    return S.IdentifyCUDAPreference(Caller, FetchDecl(Match));
>>> +  };
>>> +
>>>    // Find the best call preference among the functions in Matches.
>>> -  Sema::CUDAFunctionPreference P, BestCFP = Sema::CFP_Never;
>>> -  for (auto const &Match : Matches) {
>>> -    P = S.IdentifyCUDAPreference(Caller, FetchDecl(Match));
>>> -    if (P > BestCFP)
>>> -      BestCFP = P;
>>> -  }
>>> +  Sema::CUDAFunctionPreference BestCFP = GetCFP(*std::max_element(
>>> +      Matches.begin(), Matches.end(),
>>> +      [&](const T &M1, const T &M2) { return GetCFP(M1) < GetCFP(M2); }));
>>>
>>>    // Erase all functions with lower priority.
>>> -  for (unsigned I = 0, N = Matches.size(); I != N;)
>>> -    if (S.IdentifyCUDAPreference(Caller, FetchDecl(Matches[I])) < BestCFP) {
>>> -      Matches[I] = Matches[--N];
>>> -      Matches.resize(N);
>>> -    } else {
>>> -      ++I;
>>> -    }
>>> +  Matches.erase(llvm::remove_if(
>>> +      Matches, [&](const T &Match) { return GetCFP(Match) < BestCFP; }));
>>
>> This is the single-element form of SmallVector::erase, so it won't
>> remove ALL functions with lower priority. I sense a lack of test case
>> here.
>>
>>>  }
>>>
>>>  void Sema::EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
>>>
>>>
>>> _______________________________________________
>>> cfe-commits mailing list
>>> cfe-commits at lists.llvm.org
>>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits


More information about the cfe-commits mailing list