 
   __global__ void extractP(float* __restrict__          output,
                            const uint4* __restrict__    pyramid,
+                           const float* __restrict__    field,
                            const uint8_t* __restrict__  index_cases,
                            const uint8_t* __restrict__  index_table,
                            const size_t                 field_row_stride,

                                 const uint3                  field_size,
                                 const uint32_t               output_count,
                                 const uint2                  chunks,
                                 const float3                 scale)
  {
    if (threadIdx.x == 0) {
      uint32_t output_count_clamped = min(output_count, pyramid[8].x);
      if (output_count_clamped) {
        extractP<<<(output_count_clamped + 255) / 256, 256>>>(output,
                                                              pyramid,
                                                              index_cases,
                                                              index_table,
                                                              field_row_stride,

  if (output_buffer) {
    ::runExtractionP<<<1, 32, 0, stream>>>(reinterpret_cast<float*>(output_buffer),
                                           ctx->pyramid,
                                           field_d,
                                           ctx->index_cases_d,
                                           ctx->tables->index_table,
                                           field_row_stride,
