Disable column_view data accessors for unsupported types (#7725)

Fixes #7712

`column_view` provides data accessors like `column_view::data<T>` and `column_view::begin<T>`. These accessors are only valid for fixed-width primitive types that can be constructed by simply casting the underlying `void*` to `T*`.

However, the accessors never actually enforced this rule, e.g., `column_view::data<struct_view>` should fail to compile.

This PR disables these accessors for invalid types. 

This uncovered a number of places that were erroneously instantiating `column_view` accessors, which would lead to silent failures (e.g., `scatter` was failing silently for `struct` columns).

I added a few new things to aid me in this effort:

- `CUDF_ENABLE_IF` macro to make it easier to SFINAE.
- `is_rep_layout_compatbile<T>()` to identify types that are layout compatible with their rep (e.g., `duration_ns` is layout compatible with its `int64_t` rep. The `decimal32` type is _not_ layout compatible with it's `int32_t` rep).
- `column_device_view::has_element_accessor<T>()` identifies if `column_device_view::element<T>()` has a valid overload.

Authors:
  - Jake Hemstad (@jrhemstad)

Approvers:
  - Christopher Harris (@cwharris)
  - Conor Hoekstra (@codereport)
  - Vyas Ramasubramani (@vyasr)

URL: #7725

cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu

@@ -27,6 +27,7 @@
 
 #include <cudf/detail/utilities/cuda.cuh>
 
+#include <cudf/utilities/traits.hpp>
 #include <random>
 #include <type_traits>
 #include "../fixture/benchmark_fixture.hpp"
@@ -87,14 +88,20 @@ __global__ void host_dispatching_kernel(mutable_column_device_view source_column
 
 template <FunctorType functor_type>
 struct ColumnHandle {
-  template <typename ColumnType>
+  template <typename ColumnType, CUDF_ENABLE_IF(cudf::is_rep_layout_compatible<ColumnType>())>
   void operator()(mutable_column_device_view source_column, int work_per_thread)
   {
     cudf::detail::grid_1d grid_config{source_column.size(), block_size};
     int grid_size = grid_config.num_blocks;
     // Launch the kernel.
     host_dispatching_kernel<functor_type, ColumnType><<<grid_size, block_size>>>(source_column);
   }
+
+  template <typename ColumnType, CUDF_ENABLE_IF(not cudf::is_rep_layout_compatible<ColumnType>())>
+  void operator()(mutable_column_device_view source_column, int work_per_thread)
+  {
+    CUDF_FAIL("Invalid type to benchmark.");
+  }
 };
 
 // The following is for DEVICE_DISPATCHING:
@@ -104,12 +111,18 @@ struct ColumnHandle {
 // n_rows * n_cols.
 template <FunctorType functor_type>
 struct RowHandle {
-  template <typename T>
+  template <typename T, CUDF_ENABLE_IF(cudf::is_rep_layout_compatible<T>())>
   __device__ void operator()(mutable_column_device_view source, cudf::size_type index)
   {
     using F                 = Functor<T, functor_type>;
     source.data<T>()[index] = F::f(source.data<T>()[index]);
   }
+
+  template <typename T, CUDF_ENABLE_IF(not cudf::is_rep_layout_compatible<T>())>
+  __device__ void operator()(mutable_column_device_view source, cudf::size_type index)
+  {
+    cudf_assert(false && "Unsupported type.");
+  }
 };
 
 // This is for DEVICE_DISPATCHING

cpp/include/cudf/ast/detail/transform.cuh

@@ -20,10 +20,12 @@
 #include <cudf/ast/operators.hpp>
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/utilities/assert.cuh>
 #include <cudf/scalar/scalar_device_view.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>
+#include <cudf/utilities/traits.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -55,10 +57,19 @@ struct row_output {
    * @param row_index Row index of data column.
    * @param result Value to assign to output.
    */
-  template <typename Element>
+  template <typename Element, CUDF_ENABLE_IF(is_rep_layout_compatible<Element>())>
   __device__ void resolve_output(detail::device_data_reference device_data_reference,
                                  cudf::size_type row_index,
                                  Element result) const;
+  // Definition below after row_evaluator is a complete type
+
+  template <typename Element, CUDF_ENABLE_IF(not is_rep_layout_compatible<Element>())>
+  __device__ void resolve_output(detail::device_data_reference device_data_reference,
+                                 cudf::size_type row_index,
+                                 Element result) const
+  {
+    cudf_assert(false && "Invalid type in resolve_output.");
+  }
 
  private:
   row_evaluator const& evaluator;
@@ -167,7 +178,7 @@ struct row_evaluator {
    * @param row_index Row index of data column.
    * @return Element
    */
-  template <typename Element>
+  template <typename Element, CUDF_ENABLE_IF(column_device_view::has_element_accessor<Element>())>
   __device__ Element resolve_input(detail::device_data_reference device_data_reference,
                                    cudf::size_type row_index) const
   {
@@ -187,6 +198,15 @@ struct row_evaluator {
     }
   }
 
+  template <typename Element,
+            CUDF_ENABLE_IF(not column_device_view::has_element_accessor<Element>())>
+  __device__ Element resolve_input(detail::device_data_reference device_data_reference,
+                                   cudf::size_type row_index) const
+  {
+    cudf_assert(false && "Unsupported type in resolve_input.");
+    return {};
+  }
+
   /**
    * @brief Callable to perform a unary operation.
    *
@@ -249,7 +269,7 @@ struct row_evaluator {
   mutable_column_device_view* output_column;
 };
 
-template <typename Element>
+template <typename Element, std::enable_if_t<is_rep_layout_compatible<Element>()>*>
 __device__ void row_output::resolve_output(detail::device_data_reference device_data_reference,
                                            cudf::size_type row_index,
                                            Element result) const