Skip to content
This repository has been archived by the owner on Mar 21, 2024. It is now read-only.

thrust::reduce doesn't work with cuda::proclaim_return_type #1920

Closed
lilohuang opened this issue Apr 10, 2023 · 1 comment
Closed

thrust::reduce doesn't work with cuda::proclaim_return_type #1920

lilohuang opened this issue Apr 10, 2023 · 1 comment

Comments

@lilohuang
Copy link
Contributor

lilohuang commented Apr 10, 2023

Hi @allisonvacanti @senior-zero @jrhemstad,

In the newer version of the CUDA 12.1 SDK with the newer thrust library (> 2.0), the thrust::reduce function can emit a compile-time error when querying the return type of a device-only lambda from host code. The compiler error suggests using cuda::proclaim_return_type if you don't want to change the lambda to a host device lambda or named function object.

However, despite applying cuda::proclaim_return_type(), I still encountered compilation failure. I have provided a standalone bug reproducer at https://cuda.godbolt.org/z/ad8aTzKfT.

I am wondering if this is an expected result or if I missed something.

#include <cuda/functional>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/reduce.h>
#include <iostream>

int main() {

   thrust::device_vector<uint64_t> v(1024, 1);

#define USE_PROCLIAM_RETURN_TYPE
#ifdef USE_PROCLIAM_RETURN_TYPE
   std::cout << thrust::reduce(thrust::device,
      v.begin(), v.end(), static_cast<uint64_t>(0),
      cuda::proclaim_return_type<uint64_t>(
         [] __device__(uint64_t a, uint64_t b) -> uint64_t {
            return a + b;
         })) << std::endl;
#else
   std::cout << thrust::reduce(thrust::device,
      v.begin(), v.end(), static_cast<uint64_t>(0),
      [] __host__ __device__(uint64_t a, uint64_t b) -> uint64_t {
         return a + b;
      }) << std::endl;
#endif

   return 0;
}

Thank you so much,
Lilo

@gevtushenko
Copy link
Collaborator

Hello @lilohuang!

Your code is absolutely fine. The issue is in the initial version of the proclaim_return_type. It was addressed in the following PR. Your reproducer seems to work fine with the trunk version of libcu++.

Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.
Labels
None yet
Projects
Archived in project
Development

No branches or pull requests

2 participants