Feat/debug logger

[nathanielsimard]

• Fix cuda slicing pointer cleanup
• Fix cuda broadcast scalar assign

Example

CUBECL_DEBUG_LOG=stdout cargo run --example gelu --release --features cuda
    Finished `release` profile [optimized] target(s) in 0.11s
     Running `target/release/examples/gelu`

======== Compiled Kernel ========
name: gelu::gelu_array::GeluArray<
    cubecl_core::frontend::element::float::F32,
    cubecl_cuda::runtime::CudaRuntime,
>
cube_dim: (4, 1, 1)
shared_memory: 0 bytes
source:

typedef unsigned int uint;

extern "C" __global__ void kernel(float input_0[], float output_0[],
                                  uint info[]) {

  int3 absoluteIdx = make_int3(blockIdx.x * blockDim.x + threadIdx.x,
                               blockIdx.y * blockDim.y + threadIdx.y,
                               blockIdx.z * blockDim.z + threadIdx.z);

  uint idxGlobal =
      (absoluteIdx.z * gridDim.x * blockDim.x * gridDim.y * blockDim.y) +
      (absoluteIdx.y * gridDim.x * blockDim.x) + absoluteIdx.x;
  uint rank = info[0];
  uint rank_2 = rank * 2;
  uint l_0_0;
  bool l_0_1;
  float l_0_2;
  float l_0_3;
  l_0_0 = info[(2 * 2 * info[0]) + 1];
  l_0_1 = idxGlobal < l_0_0;
  if (l_0_1) {
    l_0_2 = input_0[idxGlobal];
    l_0_3 = sqrt(float(2.0));
    l_0_3 = l_0_2 / l_0_3;
    l_0_3 = erf(l_0_3);
    l_0_3 = l_0_3 + float(1.0);
    l_0_2 = l_0_2 * l_0_3;
    l_0_2 = l_0_2 / float(2.0);
    output_0[idxGlobal] = l_0_2;
  }
}
=================================

[louisfd]

LGTM