oneapi-src / oneapi-src/SYCLomatic

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

Open
#2,192 3 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

bug
Dominant language
LLVM
Stars
291
Forks
99
Avg merge
5d 5h
Merged PRs (30d)
1

Description

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

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with the supplied C++ reproducer and inspect SYCLomatic's handling of external template declarations during CUDA-to-SYCL migration. Compare the migrated output with the reported compiler diagnostics and icpx behavior; done means the reproducer migrates without the false-negative errors.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Stale
Clarity
Clearly specified
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.