iree/hal/local/executable_library_benchmark.c
// NOTE: error handling is here just for better diagnostics: it is not tracking
// allocations correctly and will leak. Don't use this as an example for how to
// write robust code.
error handling is here just for better diagnostics: it is not tracking
allocations correctly and will leak. Don't use this as an example for how to
write robust code.
error handling is here just for better diagnostics : it is not tracking allocations correctly and will leak . Don ' t use this as an example for how to write robust code .
static iree_status_t iree_hal_executable_library_run(
    const iree_benchmark_def_t* benchmark_def,
    iree_benchmark_state_t* benchmark_state) {
  iree_allocator_t host_allocator = benchmark_state->host_allocator;
  iree_hal_executable_loader_t* executable_loader = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_executable_library_create_loader(
      host_allocator, &executable_loader));
  iree_hal_executable_spec_t executable_spec;
  iree_hal_executable_spec_initialize(&executable_spec);
  executable_spec.caching_mode =
      IREE_HAL_EXECUTABLE_CACHING_MODE_ALLOW_OPTIMIZATION |
      IREE_HAL_EXECUTABLE_CACHING_MODE_ALIAS_PROVIDED_DATA |
      IREE_HAL_EXECUTABLE_CACHING_MODE_DISABLE_VERIFICATION;
  executable_spec.executable_format =
      iree_make_cstring_view(FLAG_executable_format);
  IREE_RETURN_IF_ERROR(iree_file_read_contents(
      FLAG_executable_file, host_allocator,
      (iree_byte_span_t*)&executable_spec.executable_data));
  executable_spec.executable_layout_count = 0;
  executable_spec.executable_layouts = NULL;
  iree_hal_executable_t* executable = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_executable_loader_try_load(
      executable_loader, &executable_spec, &executable));
  iree_hal_local_executable_t* local_executable =
      iree_hal_local_executable_cast(executable);
  iree_byte_span_t local_memory = iree_make_byte_span(NULL, 0);
  iree_host_size_t local_memory_size =
      local_executable->dispatch_attrs
          ? local_executable->dispatch_attrs[FLAG_entry_point]
                    .local_memory_pages *
                IREE_HAL_WORKGROUP_LOCAL_MEMORY_PAGE_SIZE
          : 0;
  if (local_memory_size > 0) {
    IREE_RETURN_IF_ERROR(iree_allocator_malloc(
        host_allocator, local_memory_size, (void**)&local_memory.data));
    local_memory.data_length = local_memory_size;
  }
  iree_hal_allocator_t* heap_allocator = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_allocator_create_heap(
      iree_make_cstring_view("benchmark"), host_allocator, &heap_allocator));
  iree_hal_buffer_view_t* buffer_views[IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT];
  void* binding_ptrs[IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT];
  size_t binding_lengths[IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT];
  for (iree_host_size_t i = 0; i < dispatch_params.binding_count; ++i) {
    IREE_RETURN_IF_ERROR(iree_hal_buffer_view_parse(
        dispatch_params.bindings[i], heap_allocator, &buffer_views[i]));
    iree_hal_buffer_t* buffer = iree_hal_buffer_view_buffer(buffer_views[i]);
    iree_device_size_t buffer_length =
        iree_hal_buffer_view_byte_length(buffer_views[i]);
    iree_hal_buffer_mapping_t buffer_mapping;
    IREE_RETURN_IF_ERROR(iree_hal_buffer_map_range(
        buffer, IREE_HAL_MEMORY_ACCESS_ALL, 0, buffer_length, &buffer_mapping));
    binding_ptrs[i] = buffer_mapping.contents.data;
    binding_lengths[i] = (size_t)buffer_mapping.contents.data_length;
  }
  iree_hal_executable_dispatch_state_v0_t dispatch_state = {
      .workgroup_count = {{
          .x = FLAG_workgroup_count_x,
          .y = FLAG_workgroup_count_y,
          .z = FLAG_workgroup_count_z,
      }},
      .workgroup_size = {{
          .x = FLAG_workgroup_size_x,
          .y = FLAG_workgroup_size_y,
          .z = FLAG_workgroup_size_z,
      }},
      .push_constant_count = dispatch_params.push_constant_count,
      .push_constants = &dispatch_params.push_constants[0].ui32,
      .binding_count = dispatch_params.binding_count,
      .binding_ptrs = binding_ptrs,
      .binding_lengths = binding_lengths,
      .import_thunk = NULL,  
      .imports = NULL,       
  };
  int64_t dispatch_count = 0;
  while (iree_benchmark_keep_running(benchmark_state, 1)) {
    IREE_RETURN_IF_ERROR(iree_hal_local_executable_issue_dispatch_inline(
        local_executable, FLAG_entry_point, &dispatch_state, local_memory));
    ++dispatch_count;
  }
  int64_t total_invocations =
      dispatch_count * dispatch_state.workgroup_count.x *
      dispatch_state.workgroup_count.y * dispatch_state.workgroup_count.z;
  iree_benchmark_set_items_processed(benchmark_state, total_invocations);
  for (iree_host_size_t i = 0; i < dispatch_params.binding_count; ++i) {
    iree_hal_buffer_view_release(buffer_views[i]);
  }
  iree_hal_allocator_release(heap_allocator);
  iree_allocator_free(host_allocator,
                      (void*)executable_spec.executable_data.data);
  iree_hal_executable_release(executable);
  iree_hal_executable_loader_release(executable_loader);
  return iree_ok_status();
}
static iree_status_t iree_hal_executable_library_run ( const iree_benchmark_def_t * benchmark_def , iree_benchmark_state_t * benchmark_state ) { iree_allocator_t host_allocator = benchmark_state -> host_allocator ; iree_hal_executable_loader_t * executable_loader = NULL ; IREE_RETURN_IF_ERROR ( iree_hal_executable_library_create_loader ( host_allocator , & executable_loader ) ) ; iree_hal_executable_spec_t executable_spec ; iree_hal_executable_spec_initialize ( & executable_spec ) ; executable_spec . caching_mode = IREE_HAL_EXECUTABLE_CACHING_MODE_ALLOW_OPTIMIZATION | IREE_HAL_EXECUTABLE_CACHING_MODE_ALIAS_PROVIDED_DATA | IREE_HAL_EXECUTABLE_CACHING_MODE_DISABLE_VERIFICATION ; executable_spec . executable_format = iree_make_cstring_view ( FLAG_executable_format ) ; IREE_RETURN_IF_ERROR ( iree_file_read_contents ( FLAG_executable_file , host_allocator , ( iree_byte_span_t * ) & executable_spec . executable_data ) ) ; executable_spec . executable_layout_count = 0 ; executable_spec . executable_layouts = NULL ; iree_hal_executable_t * executable = NULL ; IREE_RETURN_IF_ERROR ( iree_hal_executable_loader_try_load ( executable_loader , & executable_spec , & executable ) ) ; iree_hal_local_executable_t * local_executable = iree_hal_local_executable_cast ( executable ) ; iree_byte_span_t local_memory = iree_make_byte_span ( NULL , 0 ) ; iree_host_size_t local_memory_size = local_executable -> dispatch_attrs ? local_executable -> dispatch_attrs [ FLAG_entry_point ] . local_memory_pages * IREE_HAL_WORKGROUP_LOCAL_MEMORY_PAGE_SIZE : 0 ; if ( local_memory_size > 0 ) { IREE_RETURN_IF_ERROR ( iree_allocator_malloc ( host_allocator , local_memory_size , ( void * * ) & local_memory . data ) ) ; local_memory . data_length = local_memory_size ; } iree_hal_allocator_t * heap_allocator = NULL ; IREE_RETURN_IF_ERROR ( iree_hal_allocator_create_heap ( iree_make_cstring_view ( " " ) , host_allocator , & heap_allocator ) ) ; iree_hal_buffer_view_t * buffer_views [ IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT ] ; void * binding_ptrs [ IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT ] ; size_t binding_lengths [ IREE_HAL_LOCAL_MAX_TOTAL_BINDING_COUNT ] ; for ( iree_host_size_t i = 0 ; i < dispatch_params . binding_count ; ++ i ) { IREE_RETURN_IF_ERROR ( iree_hal_buffer_view_parse ( dispatch_params . bindings [ i ] , heap_allocator , & buffer_views [ i ] ) ) ; iree_hal_buffer_t * buffer = iree_hal_buffer_view_buffer ( buffer_views [ i ] ) ; iree_device_size_t buffer_length = iree_hal_buffer_view_byte_length ( buffer_views [ i ] ) ; iree_hal_buffer_mapping_t buffer_mapping ; IREE_RETURN_IF_ERROR ( iree_hal_buffer_map_range ( buffer , IREE_HAL_MEMORY_ACCESS_ALL , 0 , buffer_length , & buffer_mapping ) ) ; binding_ptrs [ i ] = buffer_mapping . contents . data ; binding_lengths [ i ] = ( size_t ) buffer_mapping . contents . data_length ; } iree_hal_executable_dispatch_state_v0_t dispatch_state = { . workgroup_count = { { . x = FLAG_workgroup_count_x , . y = FLAG_workgroup_count_y , . z = FLAG_workgroup_count_z , } } , . workgroup_size = { { . x = FLAG_workgroup_size_x , . y = FLAG_workgroup_size_y , . z = FLAG_workgroup_size_z , } } , . push_constant_count = dispatch_params . push_constant_count , . push_constants = & dispatch_params . push_constants [ 0 ] . ui32 , . binding_count = dispatch_params . binding_count , . binding_ptrs = binding_ptrs , . binding_lengths = binding_lengths , . import_thunk = NULL , . imports = NULL , } ; int64_t dispatch_count = 0 ; while ( iree_benchmark_keep_running ( benchmark_state , 1 ) ) { IREE_RETURN_IF_ERROR ( iree_hal_local_executable_issue_dispatch_inline ( local_executable , FLAG_entry_point , & dispatch_state , local_memory ) ) ; ++ dispatch_count ; } int64_t total_invocations = dispatch_count * dispatch_state . workgroup_count . x * dispatch_state . workgroup_count . y * dispatch_state . workgroup_count . z ; iree_benchmark_set_items_processed ( benchmark_state , total_invocations ) ; for ( iree_host_size_t i = 0 ; i < dispatch_params . binding_count ; ++ i ) { iree_hal_buffer_view_release ( buffer_views [ i ] ) ; } iree_hal_allocator_release ( heap_allocator ) ; iree_allocator_free ( host_allocator , ( void * ) executable_spec . executable_data . data ) ; iree_hal_executable_release ( executable ) ; iree_hal_executable_loader_release ( executable_loader ) ; return iree_ok_status ( ) ; }
NOTE: error handling is here just for better diagnostics: it is not tracking
allocations correctly and will leak.
NOTE : error handling is here just for better diagnostics : it is not tracking allocations correctly and will leak .
// Register the loader used to load (or find) the executable. // Setup the specification used to perform the executable load. // This information is normally used to select the appropriate loader but in // this benchmark we only have a single one. // Load the executable data. // Setup the layouts defining how each entry point is interpreted. // NOTE: we know for the embedded library loader that this is not required. // Other loaders may need it in which case it'll have to be provided. // Perform the load, which will fail if the executable cannot be loaded or // there was an issue with the layouts. // Allocate workgroup-local memory that each invocation can use. // Allocate storage for buffers and populate them. // They only need to remain valid for the duration of the invocation and all // memory accessed by the invocation will come from here. // Setup dispatch state. // not yet implemented // not yet implemented // Execute benchmark the workgroup invocation. // Note that each iteration runs through the whole grid as it's important that // we are testing the memory access patterns: if we just ran the same single // tile processing the same exact region of memory over and over we are not // testing cache effects. /*batch_count=*/ // To get a total time per invocation we set the item count to the total // invocations dispatched. That gives us both total dispatch and single // invocation times in the reporter output. // Deallocate buffers. // Unload.