From e7913a7f3913aa5ada95830c10ee3b42748955d6 Mon Sep 17 00:00:00 2001 From: Marius Brehler Date: Wed, 14 Feb 2024 12:49:22 +0100 Subject: [PATCH] Integrate IREE at openxla/iree@c121b86 (#63) --- iree_simple_embedding/CMakeLists.txt | 2 +- iree_simple_embedding/simple_embedding.c | 37 +++++++++---------- .../simple_embedding_test.mlir | 2 +- third_party/iree | 2 +- 4 files changed, 21 insertions(+), 22 deletions(-) diff --git a/iree_simple_embedding/CMakeLists.txt b/iree_simple_embedding/CMakeLists.txt index 4bac25a..8cb53a3 100644 --- a/iree_simple_embedding/CMakeLists.txt +++ b/iree_simple_embedding/CMakeLists.txt @@ -11,7 +11,7 @@ set(_TRANSLATE_TOOL_EXECUTABLE $) # Define arguments passed to iree-compile set(_ARGS) -list(APPEND _ARGS "-iree-input-type=mhlo") +list(APPEND _ARGS "-iree-input-type=stablehlo") list(APPEND _ARGS "--output-format=vm-bytecode") list(APPEND _ARGS "-iree-hal-target-backends=vmvx") # Uncomment the line below to use vulkan-spirv backend diff --git a/iree_simple_embedding/simple_embedding.c b/iree_simple_embedding/simple_embedding.c index 94272de..b5df80b 100644 --- a/iree_simple_embedding/simple_embedding.c +++ b/iree_simple_embedding/simple_embedding.c @@ -32,17 +32,17 @@ extern const iree_const_byte_span_t load_bytecode_module_data(); iree_status_t Run() { iree_vm_instance_t* instance = NULL; - IREE_RETURN_IF_ERROR( - iree_vm_instance_create(iree_allocator_system(), &instance)); + IREE_RETURN_IF_ERROR(iree_vm_instance_create( + IREE_VM_TYPE_CAPACITY_DEFAULT, iree_allocator_system(), &instance)); IREE_RETURN_IF_ERROR(iree_hal_module_register_all_types(instance)); iree_hal_device_t* device = NULL; IREE_RETURN_IF_ERROR(create_sample_device(iree_allocator_system(), &device), "create device"); iree_vm_module_t* hal_module = NULL; - IREE_RETURN_IF_ERROR( - iree_hal_module_create(instance, device, IREE_HAL_MODULE_FLAG_SYNCHRONOUS, - iree_allocator_system(), &hal_module)); + IREE_RETURN_IF_ERROR(iree_hal_module_create( + instance, /*device_count=*/1, &device, IREE_HAL_MODULE_FLAG_SYNCHRONOUS, + iree_allocator_system(), &hal_module)); // Load bytecode module from the embedded data. const iree_const_byte_span_t module_data = load_bytecode_module_data(); @@ -78,16 +78,16 @@ iree_status_t Run() { iree_hal_dim_t shape[1] = {IREE_ARRAYSIZE(kFloat4)}; iree_hal_buffer_view_t* arg0_buffer_view = NULL; iree_hal_buffer_view_t* arg1_buffer_view = NULL; - IREE_RETURN_IF_ERROR(iree_hal_buffer_view_allocate_buffer( - iree_hal_device_allocator(device), IREE_ARRAYSIZE(shape), shape, + IREE_RETURN_IF_ERROR(iree_hal_buffer_view_allocate_buffer_copy( + device, iree_hal_device_allocator(device), IREE_ARRAYSIZE(shape), shape, IREE_HAL_ELEMENT_TYPE_FLOAT_32, IREE_HAL_ENCODING_TYPE_DENSE_ROW_MAJOR, (iree_hal_buffer_params_t){ .type = IREE_HAL_MEMORY_TYPE_DEVICE_LOCAL, .usage = IREE_HAL_BUFFER_USAGE_DEFAULT, }, iree_make_const_byte_span(kFloat4, sizeof(kFloat4)), &arg0_buffer_view)); - IREE_RETURN_IF_ERROR(iree_hal_buffer_view_allocate_buffer( - iree_hal_device_allocator(device), IREE_ARRAYSIZE(shape), shape, + IREE_RETURN_IF_ERROR(iree_hal_buffer_view_allocate_buffer_copy( + device, iree_hal_device_allocator(device), IREE_ARRAYSIZE(shape), shape, IREE_HAL_ELEMENT_TYPE_FLOAT_32, IREE_HAL_ENCODING_TYPE_DENSE_ROW_MAJOR, (iree_hal_buffer_params_t){ .type = IREE_HAL_MEMORY_TYPE_DEVICE_LOCAL, @@ -97,10 +97,10 @@ iree_status_t Run() { // Setup call inputs with our buffers. iree_vm_list_t* inputs = NULL; - IREE_RETURN_IF_ERROR(iree_vm_list_create( - /*element_type=*/NULL, - /*capacity=*/2, iree_allocator_system(), &inputs), - "can't allocate input vm list"); + IREE_RETURN_IF_ERROR( + iree_vm_list_create(iree_vm_make_undefined_type_def(), + /*capacity=*/2, iree_allocator_system(), &inputs), + "can't allocate input vm list"); iree_vm_ref_t arg0_buffer_view_ref = iree_hal_buffer_view_move_ref(arg0_buffer_view); @@ -114,10 +114,10 @@ iree_status_t Run() { // Prepare outputs list to accept the results from the invocation. // The output vm list is allocated statically. iree_vm_list_t* outputs = NULL; - IREE_RETURN_IF_ERROR(iree_vm_list_create( - /*element_type=*/NULL, - /*capacity=*/1, iree_allocator_system(), &outputs), - "can't allocate output vm list"); + IREE_RETURN_IF_ERROR( + iree_vm_list_create(iree_vm_make_undefined_type_def(), + /*capacity=*/1, iree_allocator_system(), &outputs), + "can't allocate output vm list"); // Synchronously invoke the function. IREE_RETURN_IF_ERROR(iree_vm_invoke( @@ -126,8 +126,7 @@ iree_status_t Run() { // Get the result buffers from the invocation. iree_hal_buffer_view_t* ret_buffer_view = - (iree_hal_buffer_view_t*)iree_vm_list_get_ref_deref( - outputs, 0, &iree_hal_buffer_view_descriptor); + iree_vm_list_get_buffer_view_assign(outputs, 0); if (ret_buffer_view == NULL) { return iree_make_status(IREE_STATUS_NOT_FOUND, "can't find return buffer view"); diff --git a/iree_simple_embedding/simple_embedding_test.mlir b/iree_simple_embedding/simple_embedding_test.mlir index 6596219..5dd91e2 100644 --- a/iree_simple_embedding/simple_embedding_test.mlir +++ b/iree_simple_embedding/simple_embedding_test.mlir @@ -1,5 +1,5 @@ func.func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> { - %0 = "mhlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> + %0 = "stablehlo.multiply"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> return %0 : tensor<4xf32> } diff --git a/third_party/iree b/third_party/iree index 0eeae4f..c121b86 160000 --- a/third_party/iree +++ b/third_party/iree @@ -1 +1 @@ -Subproject commit 0eeae4f712cf4df4ccfd29c671413f0b486020ea +Subproject commit c121b868ee12fd7ac3d60ea2912c248a328b09a8