MuGdxy / muda

μ-Cuda, COVER THE LAST MILE OF CUDA. With features: intellisense-friendly, structured launch, automatic cuda graph generation and updating.
https://mugdxy.github.io/muda-doc/
Apache License 2.0
149 stars 7 forks source link

Question from intellisense #56

Open Da1sypetals opened 3 months ago

Da1sypetals commented 3 months ago

These days I tried to setup project using muda in clion using the toolchain in docker. However, Clion intellisense refuse to work. I found that this is because static code analysis in conditional compilation (if constexpr) choose the last branch:

// namespace muda::details
template <typename F, typename UserTag>
    MUDA_GLOBAL void parallel_for_kernel(ParallelForCallable<F> f)
    {
        if constexpr(std::is_invocable_v<F, int>)
        {
            auto tid = blockIdx.x * blockDim.x + threadIdx.x;
            auto i   = tid;
            if(i < f.count)
            {
                f.callable(i);
            }
        }
        else if constexpr(std::is_invocable_v<F, ParallelForDetails>)
        {
            ParallelForDetails details{ParallelForType::DynamicBlocks,
                                       static_cast<int>(blockIdx.x * blockDim.x
                                                        + threadIdx.x),
                                       f.count};
            if(details.i() < details.total_num())
            {
                f.callable(details);
            }
        }
        else
        {
            static_assert(always_false_v<F>, "f must be void (int) or void (ParallelForDetails)");
        }
    }

which results in a failure assertion and thus an error. But the code compiles just fine with the same toolchain setup in Clion.

In theory, the functions implemented in type_traits are host functions, which are not supported to be called in device function. I also search the documentation page related to type traits (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-type-traits) and did not find this behavior documented.

My questions:

  1. Is this behavior documented elsewhere?
  2. Is this most probably a problem of CLion itself? I also wonder what editor/IDE you to develop with Muda. (Clion just gives hundreds of errors even if my code compiles, and I am pretty sure my toolchain is correct since I can compile and debug&run with Clion GUI)

Thanks a lot in advance.

MuGdxy commented 3 months ago
  1. I didn't find anything about that before, maybe NV documented in the newer version?
  2. I use Visual Studio most of the time, so it's ok for VS's Intellisense. Can you do some workaround to remove these Intellisense errors?

For Visual Studio, if some codes compile but give Intellisense error, I will use the IDE macro to remove the code segment:

#ifndef __INTELLISENSE__
// The code causes error
#endif
MuGdxy commented 3 months ago

There is no specification for IDE Intellisence exactly. So the result of different IDE is undefined, especially coming across cuda code.

MuGdxy commented 3 months ago

About: "In theory, the functions implemented in type_traits are host functions, which are not supported to be called in device function. " I think, the type_traits is just compile time template deduction, so, device or host make no sense.

Da1sypetals commented 3 months ago

About: "In theory, the functions implemented in type_traits are host functions, which are not supported to be called in device function. " I think, the type_traits is just compile time template deduction, so, device or host make no sense.

That explains. Thanks. I used to comprehend type_traits as function (maybe a bad habit taken from python) 😭

Da1sypetals commented 3 months ago

There is no specification for IDE Intellisence exactly. So the result of different IDE is undefined, especially coming across cuda code.

That could be some internal problem of Clion then. I will probably contact jetbrains support. Thanks!