Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

getting false negative error by SYCLomatic when migrating external CUDA host device function template #2192

Open
ArberSephirotheca opened this issue Jul 23, 2024 · 3 comments
Labels
bug Something isn't working

Comments

@ArberSephirotheca
Copy link

ArberSephirotheca commented Jul 23, 2024

Describe the bug

When I try to migrate the code containing extern template

#include <cuda_runtime.h>


template <typename Outtype>
class foo1{};

template <typename Out>
 __device__ __host__ foo1<Out> foo()
{

    return foo1<Out>{};
    
}


extern template foo1<float> foo();
extern template foo1<double> foo();

int main(){
    foo1<float> foo();
}

I got the following errors:

error: explicit instantiation of 'foo' does not refer to a function template, variable template, member function, member class, or static data member
   16 | extern template foo1<float> foo();
note: candidate template ignored: target attributes do not match
    8 |  __device__ __host__ foo1<Out> foo()
 error: explicit instantiation of 'foo' does not refer to a function template, variable template, member function, member class, or static data member
   17 | extern template foo1<double> foo();
 note: candidate template ignored: target attributes do not match
    8 |  __device__ __host__ foo1<Out> foo()

code after migration:

#include <sycl/sycl.hpp>
#include <dpct/dpct.hpp>

template <typename Outtype>
class foo1{};

template <typename Out>
 foo1<Out> foo()
{

    return foo1<Out>{};
    
}


extern template foo1<float> foo();
extern template foo1<double> foo();

int main(){
    foo1<float> foo();
}

However, icpx is able to compile the migrated code without any issues.

To reproduce

#include <cuda_runtime.h>


template <typename Outtype>
class foo1{};

template <typename Out>
 __device__ __host__ foo1<Out> foo()
{

    return foo1<Out>{};
    
}


extern template foo1<float> foo();
extern template foo1<double> foo();

int main(){
    foo1<float> foo();
}

Environment

  • OS: Linux
  • Target device and vendor: Nvidia GPU
  • DPC++ version: Intel(R) oneAPI DPC++/C++ Compiler 2024.2.0 (2024.2.0.20240602)

Additional context

No response

@ArberSephirotheca ArberSephirotheca added the bug Something isn't working label Jul 23, 2024
@ArberSephirotheca ArberSephirotheca changed the title getting false negative error by SYCLomatic when migrating external CUDA function template getting false negative error by SYCLomatic when migrating external CUDA host device function template Jul 23, 2024
@tomflinda
Copy link
Contributor

@ArberSephirotheca reproduced, thanks for your report, and we will plan to fix it.

@ArberSephirotheca
Copy link
Author

ArberSephirotheca commented Jul 25, 2024

Thank you! @tomflinda

I was wondering if this error hurt the later migration process even it is false negative? I tried to eliminate this error using user define rule to replace every extern template with empty line. However, the error still persist. it seems that user define rule is applied post migration? Can we decide when to apply the user define rule? (e.g. post migration vs pre migration)

@tomflinda
Copy link
Contributor

@ArberSephirotheca, all user-defined rules are applied after the CUDA source code is parsed. Therefore, user-defined rules cannot eliminate the parsing error.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working
Projects
None yet
Development

No branches or pull requests

2 participants