Understanding C++ function pointer pass by reference in HIP

Question

HIP is the AMD GPU programming model corresponding to the NVIDIA's CUDA. I have a code snippet from HIP source code that I can't fully understand. As a reminder, the understanding of the following code snippnet doesn't require any background knowledge of HIP, but more of a question in C++ template/function pointer.

typedef int hipLaunchParm;    
template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks,
                        std::uint32_t groupMemBytes, hipStream_t stream, Args... args) 
{
    hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, 
        hipLaunchParm{}, std::move(args)...);
}

I'm confused about the following:

If F is a function pointer, why does it need to be double referenced in the argument?
How is the first template argument typename... Args useful?
hipLaunchParm is just an alias for integer, but what is the meaning of {} when it is called in the argument?

@JesperJuhl Hip is the AMD corresponding GPU programming model for CUDA https://github.com/ROCm-Developer-Tools/HIP — jer_yin, Jan 11 '19 at 17:41
Ok. Thanks. That should have been *in the question*. Don't make assumptions that people reading your question know your domain or acronyms. Make the question contain *all* relevant information. — Jesper Juhl, Jan 11 '19 at 17:42
@JesperJuhl Thanks for reminding, question updated. I should probably avoid mentioned HIP cuz the question has less to do with anything technical related with HIP but more with C++ — jer_yin, Jan 11 '19 at 17:47
@Quasimodo'sclone Sure, I added a deeplink in the description. Thanks for the Variadic_template link, that is a major gap in understanding how the code works. — jer_yin, Jan 11 '19 at 19:40

eerorika · Accepted Answer · 2019-01-11T18:09:58.407

5

If F is a function pointer, why does it need to be double referenced in the argument?

F isn't a function pointer necessarily. That is just the default type. You can pass any^† callable as long as it is invokable with the given arguments, and you want to avoid copying stateful function objects when that's not necessary. Some might not even be copyable. That may be the reason why they use a reference here.

^† as far as C++ is concerned. I don't know about restrictions that HIP / CUDA might have.

How is the first template argument typename... Args useful?

It allows passing a variable number of arguments into the delegated function.

hipLaunchParm is just a alias for integer, but what is the meaning of {} when it is called in the argument?

T{} is a syntax for value initialisation of a temporary. In case of integer, this means that zero is passed as an argument.

edited Jan 11 '19 at 18:09

answered Jan 11 '19 at 17:56

eerorika

232,697
12
197
326

Now it looks like my gap is in understanding of c++11 features. The first one is about more generic nature of rvalue reference. Second one is variadic template. Third one is initializer list. – jer_yin Jan 11 '19 at 22:31
@jer_yin as far as the rvalue reference (technically forwarding reference, since the type is deduced) is concerned: I have no idea why they use that instead of a const lvalue reference. – eerorika Jan 11 '19 at 22:32
I find a good resource [here](https://stackoverflow.com/questions/5481539/what-does-t-double-ampersand-mean-in-c11), that explains the benefit of rvalue reference versus const lvalue. – jer_yin Jan 11 '19 at 22:35
@jer_yin I know exactly what rvalue references are used for, but the code here shows no need to use it for the function argument. – eerorika Jan 11 '19 at 22:36

Understanding C++ function pointer pass by reference in HIP

1 Answers1