Remove util_device.cuh from iterator headers to enable online compilation

[leofang]

Description

Related to #403.

When using NVRTC (+Jitify) to compile these two headers

#include <cub/block/block_reduce.cuh>
#include <cub/block/block_load.cuh>

util_device.cuh ended up being included by one of the CUB headers, and this function

cccl/cub/cub/util_device.cuh

Lines 185 to 192 in eea703e

	// Host code, only safely usable in C++11 or newer, where thread-safe
	// initialization of static locals is guaranteed. This is a separate function
	// to avoid defining a local static in a host/device function.
	__host__ inline int DeviceCountCachedValue()
	{
	static ValueCache<int, DeviceCountUncached> cache;
	return cache.value;
	}

causes the following compiler error

../util_device.cuh(191): error: dynamic initialization is not supported for a function-scope static __device__ variable within a __device__/__global__ function

Upon a closer look, I realized that by commenting out DeviceCountCachedValue, the struct PerDeviceAttributeCache, and everything depending on them (in the same file, so the radius of blast is small), I was able to make NVRTC happy. Basically, hiding all __host__ only functions and structs from NVRTC should suffice.

@jrhemstad pointed out offline that this header doesn't seem to be used in cache_modified_input_iterator.cuh, and I checked locally that it's true for all iterator headers. Hence this patch.

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[jrhemstad]

/ok to test

[leofang]

Findings:

• Several of the device tuning/dispatch headers need ChainedPolicy, so I made those include util_device.cuh directly, as device algorithms contain host API calls and cannot be compiled by NVRTC anyway.
• In other places, util_device.cuh was included only to access AliasTemporaries. Since it's __host__ __device__ and compilable by NVRTC, I moved it to a separate util header and made it replace the include of util_device.cuh

[leofang]

Can anyone help kick off CI to validate the changes?

[miscco]

/ok to test

[leofang]

Thanks, Michael, it seems the tests are green 🙂

[miscco]

Minor nits

[gevtushenko]

Thank you for cleaning up the headers! It's much better now. A minor comment below.

[miscco]

/ok to test