Add samples to build/run vector multiplication with RVV instructions Add both int32 and float examples. mhlo.dot MLIR (i8,i8)->i32 (int8 dot product to int32 out) doesn't support RVV codegen at clang trunk. The binary can run successfully on qemu-system. renode will hit illegal instuction for the float example. Change-Id: I6c87211542fea485eda4d3de7da8250ec60e2a01

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 62a54f3..4b7e783 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt

@@ -37,4 +37,15 @@
   add_definitions(-DIREE_USER_CONFIG_H="${SPRINGBOK_CONFIG_HEADER}")
 endif()
 
+# Build IREE runtime libraries.
 add_subdirectory($ENV{ROOTDIR}/toolchain/iree iree)
+
+# Apply IREE's CMake variables and build options so we can use IREE build
+# functions properly in this project.
+include($ENV{ROOTDIR}/toolchain/iree/build_tools/cmake/iree_copts.cmake)
+
+# Include springbok log functions.
+include_directories(BEFORE SYSTEM "$ENV{ROOTDIR}/sw/vec/springbok/include")
+
+# Add the included directory here.
+add_subdirectory(samples)

diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt
new file mode 100644
index 0000000..0e9c88b
--- /dev/null
+++ b/samples/CMakeLists.txt

@@ -0,0 +1 @@
+iree_add_all_subdirs()

diff --git a/samples/simple_vec_mul/CMakeLists.txt b/samples/simple_vec_mul/CMakeLists.txt
new file mode 100644
index 0000000..647effc
--- /dev/null
+++ b/samples/simple_vec_mul/CMakeLists.txt

@@ -0,0 +1,102 @@
+#-------------------------------------------------------------------------------
+# Build the mlir bytecode modules with iree-translate. Note the last two flags
+# and "+experimental-v" are for RVV support.
+# https://github.com/llvm/llvm-project/blob/f4ec30d/llvm/lib/Target/RISCV/RISCVSubtarget.cpp#L30-L46
+#-------------------------------------------------------------------------------
+
+iree_bytecode_module(
+  NAME
+    simple_float_mul_bytecode_module_dylib
+  SRC
+    "simple_float_mul.mlir"
+  C_IDENTIFIER
+    "samples_simple_vec_mul_simple_float_mul_bytecode_module_dylib"
+  FLAGS
+    "-iree-input-type=mhlo"
+    "-iree-mlir-to-vm-bytecode-module"
+    "-iree-hal-target-backends=dylib-llvm-aot"
+    "-iree-llvm-target-triple=riscv32-pc-linux-elf"
+    "-iree-llvm-target-cpu=generic-rv32"
+    "-iree-llvm-target-cpu-features=+m,+f,+experimental-v"
+    "-iree-llvm-target-abi=ilp32"
+    "-iree-llvm-link-embedded=true"
+    "-iree-llvm-debug-symbols=false"
+    "-riscv-v-vector-bits-min=512"
+    "-riscv-v-fixed-length-vector-lmul-max=8"
+  PUBLIC
+)
+
+# Temporarily disable RVV for the int32 version so the binary can be run on
+# renode.
+# TODO(b/195201585): Re-enable RVV.
+iree_bytecode_module(
+  NAME
+    simple_int_mul_bytecode_module_dylib
+  SRC
+    "simple_int_mul.mlir"
+  C_IDENTIFIER
+    "samples_simple_vec_mul_simple_int_mul_bytecode_module_dylib"
+  FLAGS
+    "-iree-input-type=mhlo"
+    "-iree-mlir-to-vm-bytecode-module"
+    "-iree-hal-target-backends=dylib-llvm-aot"
+    "-iree-llvm-target-triple=riscv32-pc-linux-elf"
+    "-iree-llvm-target-cpu=generic-rv32"
+    "-iree-llvm-target-cpu-features=+m,+f,+experimental-v"
+    "-iree-llvm-target-abi=ilp32"
+    "-iree-llvm-link-embedded=true"
+    "-iree-llvm-debug-symbols=false"
+    "-riscv-v-vector-bits-min=512"
+    "-riscv-v-fixed-length-vector-lmul-max=8"
+  PUBLIC
+)
+
+#-------------------------------------------------------------------------------
+# Binaries to execute the MLIR bytecode modules
+#-------------------------------------------------------------------------------
+
+
+# If the program requires a larger stack size, add
+#
+#   LINKOPTS
+#     "LINKER:--defsym=__stack_size__=<new stack size>"
+#
+# to increase it.
+
+iree_cc_binary(
+  NAME
+    simple_float_vec_mul_embedded_sync
+  SRCS
+    "device_embedded_sync.c"
+    "float_vec.c"
+    "simple_vec_mul.c"
+  DEPS
+    ::simple_float_mul_bytecode_module_dylib_c
+    iree::base
+    iree::hal
+    iree::hal::local
+    iree::hal::local::loaders::embedded_library_loader
+    iree::hal::local::sync_driver
+    iree::modules::hal
+    iree::vm
+    iree::vm::bytecode_module
+)
+
+iree_cc_binary(
+  NAME
+    simple_int_vec_mul_embedded_sync
+  SRCS
+    "device_embedded_sync.c"
+    "int_vec.c"
+    "simple_vec_mul.c"
+  DEPS
+    ::simple_int_mul_bytecode_module_dylib_c
+    iree::base
+    iree::hal
+    iree::hal::local
+    iree::hal::local::loaders::embedded_library_loader
+    iree::hal::local::sync_driver
+    iree::modules::hal
+    iree::vm
+    iree::vm::bytecode_module
+)

diff --git a/samples/simple_vec_mul/device_embedded_sync.c b/samples/simple_vec_mul/device_embedded_sync.c
new file mode 100644
index 0000000..2f10e6b
--- /dev/null
+++ b/samples/simple_vec_mul/device_embedded_sync.c

@@ -0,0 +1,29 @@
+// An example of setting up the embedded-sync driver.
+
+#include <stddef.h>
+
+#include "iree/base/api.h"
+#include "iree/hal/api.h"
+#include "iree/hal/local/executable_loader.h"
+#include "iree/hal/local/loaders/embedded_library_loader.h"
+#include "iree/hal/local/sync_device.h"
+
+iree_status_t create_sample_device(iree_hal_device_t **device) {
+  // Set parameters for the device created in the next step.
+  iree_hal_sync_device_params_t params;
+  iree_hal_sync_device_params_initialize(&params);
+
+  iree_hal_executable_loader_t *loader = NULL;
+  IREE_RETURN_IF_ERROR(iree_hal_embedded_library_loader_create(
+      iree_hal_executable_import_provider_null(), iree_allocator_system(),
+      &loader));
+
+  iree_string_view_t identifier = iree_make_cstring_view("dylib");
+
+  // Create the synchronous device and release the loader afterwards.
+  iree_status_t status =
+      iree_hal_sync_device_create(identifier, &params, /*loader_count=*/1,
+                                  &loader, iree_allocator_system(), device);
+  iree_hal_executable_loader_release(loader);
+  return status;
+}

diff --git a/samples/simple_vec_mul/float_vec.c b/samples/simple_vec_mul/float_vec.c
new file mode 100644
index 0000000..eaa41e1
--- /dev/null
+++ b/samples/simple_vec_mul/float_vec.c

@@ -0,0 +1,77 @@
+// Float simple_mul bytecode loading and input/output processes
+
+#include "iree/base/api.h"
+#include "iree/hal/api.h"
+
+// Compiled module embedded here to avoid file IO:
+#include "samples/simple_vec_mul/simple_float_mul_bytecode_module_dylib_c.h"
+
+const iree_const_byte_span_t load_bytecode_module_data() {
+  const struct iree_file_toc_t *module_file_toc =
+      samples_simple_vec_mul_simple_float_mul_bytecode_module_dylib_create();
+  return iree_make_const_byte_span(module_file_toc->data,
+                                   module_file_toc->size);
+}
+
+iree_status_t prepare_input_hal_buffer_views(
+    iree_hal_device_t *device, const int buffer_length, void **arg0_buffer,
+    void **arg1_buffer, iree_hal_buffer_view_t **arg0_buffer_view,
+    iree_hal_buffer_view_t **arg1_buffer_view) {
+  iree_status_t result = iree_ok_status();
+  *arg0_buffer =
+      iree_aligned_alloc(sizeof(uint32_t), sizeof(float) * buffer_length);
+  if (*arg0_buffer == NULL) {
+    result = iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED);
+  }
+  if (iree_status_is_ok(result)) {
+    *arg1_buffer =
+        iree_aligned_alloc(sizeof(uint32_t), sizeof(float) * buffer_length);
+    if (*arg1_buffer == NULL) {
+      result = iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED);
+    }
+  }
+
+  // Populate initial values
+  // arg0 = 0, 1/4, 1/2, 3/4... 1023/4
+  // arg1 = 0, 1/2, 1, 3/2... 1023/2
+  for (int i = 0; i < buffer_length; ++i) {
+    ((float *)*arg0_buffer)[i] = i / 4.0f;
+    ((float *)*arg1_buffer)[i] = i / 2.0f;
+  }
+  // Wrap buffers in shaped buffer views.
+  // The buffers can be mapped on the CPU and that can also be used
+  // on the device. Not all devices support this, but the ones we have now do.
+  iree_hal_dim_t shape[1] = {buffer_length};
+
+  iree_hal_memory_type_t input_memory_type =
+      IREE_HAL_MEMORY_TYPE_HOST_LOCAL | IREE_HAL_MEMORY_TYPE_DEVICE_VISIBLE;
+
+  result = iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+      iree_hal_device_allocator(device), shape, IREE_ARRAYSIZE(shape),
+      IREE_HAL_ELEMENT_TYPE_FLOAT_32, input_memory_type,
+      IREE_HAL_MEMORY_ACCESS_READ, IREE_HAL_BUFFER_USAGE_ALL,
+      iree_make_byte_span(*arg0_buffer, sizeof(int32_t) * buffer_length),
+      iree_allocator_null(), arg0_buffer_view);
+
+  if (iree_status_is_ok(result)) {
+    result = iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+        iree_hal_device_allocator(device), shape, IREE_ARRAYSIZE(shape),
+        IREE_HAL_ELEMENT_TYPE_FLOAT_32, input_memory_type,
+        IREE_HAL_MEMORY_ACCESS_READ, IREE_HAL_BUFFER_USAGE_ALL,
+        iree_make_byte_span(*arg1_buffer, sizeof(int32_t) * buffer_length),
+        iree_allocator_null(), arg1_buffer_view);
+  }
+  return result;
+}
+
+iree_status_t check_output_data(iree_hal_buffer_mapping_t *mapped_memory) {
+  iree_status_t result = iree_ok_status();
+  for (int i = 0; i < mapped_memory->contents.data_length / sizeof(float);
+       ++i) {
+    if (((const float *)mapped_memory->contents.data)[i] != i * i / 8.0f) {
+      result = iree_make_status(IREE_STATUS_UNKNOWN, "result mismatches");
+      break;
+    }
+  }
+  return result;
+}

diff --git a/samples/simple_vec_mul/int_vec.c b/samples/simple_vec_mul/int_vec.c
new file mode 100644
index 0000000..946aa63
--- /dev/null
+++ b/samples/simple_vec_mul/int_vec.c

@@ -0,0 +1,77 @@
+// Integer simple_mul bytecode loading and input/output processes
+
+#include "iree/base/api.h"
+#include "iree/hal/api.h"
+
+// Compiled module embedded here to avoid file IO:
+#include "samples/simple_vec_mul/simple_int_mul_bytecode_module_dylib_c.h"
+
+const iree_const_byte_span_t load_bytecode_module_data() {
+  const struct iree_file_toc_t *module_file_toc =
+      samples_simple_vec_mul_simple_int_mul_bytecode_module_dylib_create();
+  return iree_make_const_byte_span(module_file_toc->data,
+                                   module_file_toc->size);
+}
+
+iree_status_t prepare_input_hal_buffer_views(
+    iree_hal_device_t *device, const int buffer_length, void **arg0_buffer,
+    void **arg1_buffer, iree_hal_buffer_view_t **arg0_buffer_view,
+    iree_hal_buffer_view_t **arg1_buffer_view) {
+  iree_status_t result = iree_ok_status();
+  *arg0_buffer =
+      iree_aligned_alloc(sizeof(uint32_t), sizeof(int32_t) * buffer_length);
+  if (*arg0_buffer == NULL) {
+    result = iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED);
+  }
+  if (iree_status_is_ok(result)) {
+    *arg1_buffer =
+        iree_aligned_alloc(sizeof(uint32_t), sizeof(int32_t) * buffer_length);
+    if (*arg1_buffer == NULL) {
+      result = iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED);
+    }
+  }
+
+  // Populate initial values
+  // arg0 = 0, 0, 1, 1,..., 511
+  // arg1 = 0, 1, 2, 3,..., 1023
+  for (int i = 0; i < buffer_length; ++i) {
+    ((int32_t *)*arg0_buffer)[i] = i >> 1;
+    ((int32_t *)*arg1_buffer)[i] = i;
+  }
+  // Wrap buffers in shaped buffer views.
+  // The buffers can be mapped on the CPU and that can also be used
+  // on the device. Not all devices support this, but the ones we have now do.
+  iree_hal_dim_t shape[1] = {buffer_length};
+
+  iree_hal_memory_type_t input_memory_type =
+      IREE_HAL_MEMORY_TYPE_HOST_LOCAL | IREE_HAL_MEMORY_TYPE_DEVICE_VISIBLE;
+
+  result = iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+      iree_hal_device_allocator(device), shape, IREE_ARRAYSIZE(shape),
+      IREE_HAL_ELEMENT_TYPE_SINT_32, input_memory_type,
+      IREE_HAL_MEMORY_ACCESS_READ, IREE_HAL_BUFFER_USAGE_ALL,
+      iree_make_byte_span(*arg0_buffer, sizeof(int32_t) * buffer_length),
+      iree_allocator_null(), arg0_buffer_view);
+
+  if (iree_status_is_ok(result)) {
+    result = iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+        iree_hal_device_allocator(device), shape, IREE_ARRAYSIZE(shape),
+        IREE_HAL_ELEMENT_TYPE_SINT_32, input_memory_type,
+        IREE_HAL_MEMORY_ACCESS_READ, IREE_HAL_BUFFER_USAGE_ALL,
+        iree_make_byte_span(*arg1_buffer, sizeof(int32_t) * buffer_length),
+        iree_allocator_null(), arg1_buffer_view);
+  }
+  return result;
+}
+
+iree_status_t check_output_data(iree_hal_buffer_mapping_t *mapped_memory) {
+  iree_status_t result = iree_ok_status();
+  for (int i = 0; i < mapped_memory->contents.data_length / sizeof(int32_t);
+       ++i) {
+    if (((const int32_t *)mapped_memory->contents.data)[i] != (i >> 1) * i) {
+      result = iree_make_status(IREE_STATUS_UNKNOWN, "result mismatches");
+      break;
+    }
+  }
+  return result;
+}

diff --git a/samples/simple_vec_mul/simple_float_mul.mlir b/samples/simple_vec_mul/simple_float_mul.mlir
new file mode 100644
index 0000000..7a928ff
--- /dev/null
+++ b/samples/simple_vec_mul/simple_float_mul.mlir

@@ -0,0 +1,5 @@
+func @simple_mul(%arg0: tensor<1024xf32>, %arg1: tensor<1024xf32>) -> tensor<1024xf32>
+{
+  %0 = "mhlo.multiply"(%arg0, %arg1) : (tensor<1024xf32>, tensor<1024xf32>) -> tensor<1024xf32>
+  return %0 : tensor<1024xf32>
+}

diff --git a/samples/simple_vec_mul/simple_int_mul.mlir b/samples/simple_vec_mul/simple_int_mul.mlir
new file mode 100644
index 0000000..9aa9a22
--- /dev/null
+++ b/samples/simple_vec_mul/simple_int_mul.mlir

@@ -0,0 +1,5 @@
+func @simple_mul(%arg0: tensor<1024xi32>, %arg1: tensor<1024xi32>) -> tensor<1024xi32>
+{
+  %0 = "mhlo.multiply"(%arg0, %arg1) : (tensor<1024xi32>, tensor<1024xi32>) -> tensor<1024xi32>
+  return %0 : tensor<1024xi32>
+}

diff --git a/samples/simple_vec_mul/simple_vec_mul.c b/samples/simple_vec_mul/simple_vec_mul.c
new file mode 100644
index 0000000..a2951ea
--- /dev/null
+++ b/samples/simple_vec_mul/simple_vec_mul.c

@@ -0,0 +1,169 @@
+// An example based on iree/samples/simple_embedding. Test a 1024-element int32
+// multiplication with vector extension ISA.
+
+#include <springbok.h>
+#include <stdio.h>
+
+#include "iree/base/api.h"
+#include "iree/hal/api.h"
+#include "iree/modules/hal/module.h"
+#include "iree/vm/api.h"
+#include "iree/vm/bytecode_module.h"
+
+// A function to create the HAL device from the different backend targets.
+// The HAL device is returned based on the implementation, and it must be
+// released by the caller.
+extern iree_status_t create_sample_device(iree_hal_device_t **device);
+
+extern const iree_const_byte_span_t load_bytecode_module_data();
+
+// Prepare the input buffers and buffer_views based on the data type. They must
+// be released by the caller.
+extern iree_status_t prepare_input_hal_buffer_views(
+    iree_hal_device_t *device, const int buffer_length, void **arg0_buffer,
+    void **arg1_buffer, iree_hal_buffer_view_t **arg0_buffer_view,
+    iree_hal_buffer_view_t **arg1_buffer_view);
+
+extern iree_status_t check_output_data(
+    iree_hal_buffer_mapping_t *mapped_memory);
+
+iree_status_t run() {
+  IREE_RETURN_IF_ERROR(iree_hal_module_register_types());
+
+  iree_vm_instance_t *instance = NULL;
+  iree_status_t result =
+      iree_vm_instance_create(iree_allocator_system(), &instance);
+
+  iree_hal_device_t *device = NULL;
+  if (iree_status_is_ok(result)) {
+    result = create_sample_device(&device);
+  }
+  iree_vm_module_t *hal_module = NULL;
+  if (iree_status_is_ok(result)) {
+    result =
+        iree_hal_module_create(device, iree_allocator_system(), &hal_module);
+  }
+  // Load bytecode module from the embedded data.
+  const iree_const_byte_span_t module_data = load_bytecode_module_data();
+
+  iree_vm_module_t *bytecode_module = NULL;
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_bytecode_module_create(module_data, iree_allocator_null(),
+                                            iree_allocator_system(),
+                                            &bytecode_module);
+  }
+
+  // Allocate a context that will hold the module state across invocations.
+  iree_vm_context_t *context = NULL;
+  iree_vm_module_t *modules[] = {hal_module, bytecode_module};
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_context_create_with_modules(
+        instance, &modules[0], IREE_ARRAYSIZE(modules), iree_allocator_system(),
+        &context);
+  }
+  iree_vm_module_release(hal_module);
+  iree_vm_module_release(bytecode_module);
+
+  // Lookup the entry point function.
+  // Note that we use the synchronous variant which operates on pure type/shape
+  // erased buffers.
+  const char kMainFunctionName[] = "module.simple_mul";
+  iree_vm_function_t main_function;
+  if (iree_status_is_ok(result)) {
+    result = (iree_vm_context_resolve_function(
+        context, iree_make_cstring_view(kMainFunctionName), &main_function));
+  }
+
+  // Prepare the input buffers.
+  void *arg0_buffer = NULL;
+  void *arg1_buffer = NULL;
+  iree_hal_buffer_view_t *arg0_buffer_view = NULL;
+  iree_hal_buffer_view_t *arg1_buffer_view = NULL;
+  const int kElementCount = 1024;
+  if (iree_status_is_ok(result)) {
+    result = prepare_input_hal_buffer_views(device, kElementCount, &arg0_buffer,
+                                            &arg1_buffer, &arg0_buffer_view,
+                                            &arg1_buffer_view);
+  }
+
+  // Setup call inputs with our buffers.
+  iree_vm_list_t *inputs = NULL;
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_list_create(
+        /*element_type=*/NULL,
+        /*capacity=*/2, iree_allocator_system(), &inputs);
+  }
+  iree_vm_ref_t arg0_buffer_view_ref =
+      iree_hal_buffer_view_move_ref(arg0_buffer_view);
+  iree_vm_ref_t arg1_buffer_view_ref =
+      iree_hal_buffer_view_move_ref(arg1_buffer_view);
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_list_push_ref_move(inputs, &arg0_buffer_view_ref);
+  }
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_list_push_ref_move(inputs, &arg1_buffer_view_ref);
+  }
+
+  // Prepare outputs list to accept the results from the invocation.
+  // The output vm list is allocated statically.
+  iree_vm_list_t *outputs = NULL;
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_list_create(
+        /*element_type=*/NULL,
+        /*capacity=*/1, iree_allocator_system(), &outputs);
+  }
+
+  // Invoke the function.
+  if (iree_status_is_ok(result)) {
+    result = iree_vm_invoke(context, main_function,
+                            /*policy=*/NULL, inputs, outputs,
+                            iree_allocator_system());
+  }
+
+  // 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_get_descriptor());
+  if (ret_buffer_view == NULL) {
+    result = iree_make_status(IREE_STATUS_NOT_FOUND,
+                              "can't find return buffer view");
+  }
+
+  // Read back the results and ensure we got the right values.
+  iree_hal_buffer_mapping_t mapped_memory;
+  if (iree_status_is_ok(result)) {
+    result = iree_hal_buffer_map_range(
+        iree_hal_buffer_view_buffer(ret_buffer_view),
+        IREE_HAL_MEMORY_ACCESS_READ, 0, IREE_WHOLE_BUFFER, &mapped_memory);
+  }
+  if (iree_status_is_ok(result)) {
+    result = check_output_data(&mapped_memory);
+  }
+  iree_hal_buffer_unmap_range(&mapped_memory);
+
+  iree_vm_list_release(inputs);
+  iree_vm_list_release(outputs);
+  iree_aligned_free(arg0_buffer);
+  iree_aligned_free(arg1_buffer);
+  iree_hal_device_release(device);
+  iree_vm_context_release(context);
+  iree_vm_instance_release(instance);
+  return result;
+}
+
+int main() {
+  const iree_status_t result = run();
+  int ret = (int)iree_status_code(result);
+  if (!iree_status_is_ok(result)) {
+    char *status_buffer = NULL;
+    iree_host_size_t status_buffer_length = 0;
+    iree_status_to_string(result, &status_buffer, &status_buffer_length);
+    LOG_ERROR("%.*s", (int)status_buffer_length, status_buffer);
+    free(status_buffer);
+    iree_status_free(result);
+  } else {
+    LOG_INFO("simple_vec_mul finished successfully");
+  }
+
+  return ret;
+}