[FEA] Refactor grid_1d to allow usage in JITIFY kernels

[lamarrr]

Is your feature request related to a problem? Please describe.
We have to write the global thread id calculation in many places within CUDF, specifically within JITIFY kernels.

Examples:

cudf/cpp/src/binaryop/jit/kernel.cu

Line 54 in 8e4408c

auto const start = threadIdx.x + static_cast<cudf::thread_index_type>(blockIdx.x) * blockDim.x;

cudf/cpp/src/binaryop/jit/kernel.cu

Line 73 in 8e4408c

auto const start = threadIdx.x + static_cast<cudf::thread_index_type>(blockIdx.x) * blockDim.x;

cudf/cpp/src/rolling/jit/kernel.cu

Line 54 in 8e4408c

cudf::thread_index_type i = blockIdx.x * blockDim.x + threadIdx.x;

cudf/cpp/src/transform/jit/kernel.cu

Line 53 in 8e4408c

auto const block_size = static_cast<thread_index_type>(blockDim.x);

This functionality already exists in the grid_1d class .

cudf/cpp/include/cudf/detail/utilities/cuda.cuh

Line 40 in 8e4408c

class grid_1d {

We can't use it in JITIFY kernels because the headers it is accompanied by contain host code which is problematic for JITIFY.

Some of the problematic headers are:

• cudf/utilities/default_stream.hpp
• cudf/utilities/error.hpp
• rmm/cuda_stream_view.hpp
• cub/cub.cuh

Describe the solution you'd like

These headers are not required for it to function. We should refactor out the class into a separate device-only header.