Is there a way to statically check kernel arg types for their safety

[pauleonix]

I am working on a existing CUDA C++ codebase that combines Thrust fancy iterators with eyalroz/cuda-api-wrappers. In recent versions these wrappers have started statically checking the (decayed) type of kernel arguments passed through it's API for std::is_trivially_copyable or even more recently std::is_trivially_copy_constructible. Neither of these requirements is met by Thrust fancy iterators like even the most basic thrust::counting_iterator<int>. As the two of us couldn't reach a conclusion or our own in eyalroz/cuda-api-wrappers#551, we decided to ask here if there is another option to check for correctness without compromising composability with Thrust fancy iterators.

[miscco]

Thanks a lot, that is a very good question.

Will discuss this in the coming week in the internal meeting. I need to have a look at our fancy iterator implementation though, because it is a bit dated

[eyalroz]

On my part, I definitely want to help resolve this, so please feel free to contact me if you want to discuss ideas which involve changes on both "sides", i.e. in my wrappers as well. My GitHub profile has my email, for coordinating something like that.

[pauleonix]

The relevant section of the Programming Guide says:

When a __global__ function is launched from host code, each argument type is allowed to be non-trivially copyable or non-trivially-destructible, but the processing for such types does not follow the standard C++ model, as described below. User code must ensure that this workflow does not affect program correctness.

[eyalroz]

@pauleonix : Perhaps you want to make those iterators copy-constructible?

-- snip --

I was about to say something unkind here, but decided not to.

[pauleonix]

@jrhemstad I honestly couldn't figure out what stops them from being trivially constructible last time I looked. Any pointers?

[miscco]

They are not trivially copy constructible, because they have user defined copy constructors.

The issue is that we do not have the capacity to clean up all of thrust after dropping pre C++11 support.

[eyalroz]

after dropping pre C++11 support.

isn't it easier to do cleanup after dropping legacy support?

Anyway, if the ctor's are effectively trivial, perhaps we could discuss a trait (within a detail namespace) for being either really-trivially-copy-constructible, or effectivelly-trivially-copy-constructible, i.e. so that a copy ctor may exist, but could be ignored when copying to the GPU in favor of trivial copy-construction. Then you would specialize that trait for relevant thrust ctors, and we'll be set.

[miscco]

Turns out that most of it was indeed relatively easy, so I opened a PR for it.

The issue is not whether it is easy to do something, but whether we actually have time for it. Today was a super fractured day for me, so I could take a bit of time to do something like this, but usually that is not the case.

[eyalroz]

things might get messed up if and only if the copy-constructor has side effects or somehow else relies on being called a certain number of times.

That's what I meant by being not-trivially-copy-constructible "in fact, rather than just in principle"...

[eyalroz]

ping.

[miscco]

I have opened a PR that should address most of the fancy iterators bar transform_iterator #1368

[eyalroz]

Wait, aren't you that journalist who testified before congress a while ago?

[miscco]

Oh boy, you know how many right wing bots I get on twitter for having a similar name?