diff --git a/programming_examples/basic/vector_reduce_min/run_vck5000.lit b/programming_examples/basic/vector_reduce_min/run_vck5000.lit deleted file mode 100644 index d314eea2a4..0000000000 --- a/programming_examples/basic/vector_reduce_min/run_vck5000.lit +++ /dev/null @@ -1,9 +0,0 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// REQUIRES: hsa -// -// RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir -// RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - diff --git a/programming_examples/basic/vector_reduce_min/test_vck5000.cpp b/programming_examples/basic/vector_reduce_min/test_vck5000.cpp deleted file mode 100644 index d143c329b7..0000000000 --- a/programming_examples/basic/vector_reduce_min/test_vck5000.cpp +++ /dev/null @@ -1,141 +0,0 @@ -//===- test_vck5000.cpp -----------------------------------000---*- C++ -*-===// -// -// This file is licensed under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -// Copyright (C) 2024, Advanced Micro Devices, Inc. -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "memory_allocator.h" -#include "test_library.h" - -#include "aie_data_movement.cpp" -#include "aie_inc.cpp" - -#include "hsa/hsa.h" -#include "hsa/hsa_ext_amd.h" - -constexpr int DMA_COUNT = 64; - -void hsa_check_status(const std::string func_name, hsa_status_t status) { - if (status != HSA_STATUS_SUCCESS) { - const char *status_string(new char[1024]); - hsa_status_string(status, &status_string); - std::cout << func_name << " failed: " << status_string << std::endl; - delete[] status_string; - } else { - std::cout << func_name << " success" << std::endl; - } -} - -int main(int argc, char *argv[]) { - uint64_t row = 0; - uint64_t col = 6; - - std::vector queues; - uint32_t aie_max_queue_size(0); - - aie_libxaie_ctx_t *xaie = mlir_aie_init_libxaie(); - - // This is going to initialize HSA, create a queue - // and get an agent - int ret = mlir_aie_init_device(xaie); - - if (ret) { - std::cout << "[ERROR] Error when calling mlir_aie_init_device)" - << std::endl; - return -1; - } - - // Getting access to all of the HSA agents - std::vector agents = xaie->agents; - - if (agents.empty()) { - std::cout << "No agents found. Exiting." << std::endl; - return -1; - } - - std::cout << "Found " << agents.size() << " agents" << std::endl; - - hsa_queue_t *q = xaie->cmd_queue; - - // Adding to our vector of queues - queues.push_back(q); - assert(queues.size() > 0 && "No queues were sucesfully created!"); - - mlir_aie_configure_cores(xaie); - mlir_aie_configure_switchboxes(xaie); - mlir_aie_initialize_locks(xaie); - mlir_aie_configure_dmas(xaie); - mlir_aie_start_cores(xaie); - - // Allocating some device memory - ext_mem_model_t buf0, buf1, buf2; - uint32_t *in_a = (uint32_t *)mlir_aie_mem_alloc(xaie, buf0, DMA_COUNT); - uint32_t *in_b = (uint32_t *)mlir_aie_mem_alloc(xaie, buf1, DMA_COUNT); - uint32_t *out = (uint32_t *)mlir_aie_mem_alloc( - xaie, buf2, 4 /* For some reason can't do 1 */); - mlir_aie_sync_mem_dev(buf0); - mlir_aie_sync_mem_dev(buf1); - mlir_aie_sync_mem_dev(buf2); - - if (in_a == nullptr || in_b == nullptr || out == nullptr) { - std::cout << "Could not allocate in device memory" << std::endl; - return -1; - } - - out[0] = 0xdeface; - for (int i = 0; i < DMA_COUNT; i++) { - in_a[i] = i + 100; - } - - in_a[DMA_COUNT / 2] = 3; - in_a[DMA_COUNT - 1] = 100; - - // Pass arguments in the order of dma_memcpys in the mlir - invoke_data_movement(queues[0], &agents[0], out, in_a); - - int errors = 0; - - uint32_t min_val = 1000; - for (int i = 0; i < DMA_COUNT; i++) { - uint32_t s = in_a[i]; - if (min_val > s) { - min_val = s; - } - } - - if (*out != min_val) { - errors++; - printf("[ERROR] Min value is %d but kernel returned %d\n", min_val, *out); - } - - // destroying the queue - hsa_queue_destroy(queues[0]); - - // Shutdown AIR and HSA - mlir_aie_deinit_libxaie(xaie); - - if (!errors) { - printf("PASS!\n"); - return 0; - } else { - printf("fail %d/%d.\n", errors, 1); - return -1; - } -} diff --git a/programming_examples/basic/vector_scalar_mul/run_vck5000.lit b/programming_examples/basic/vector_scalar_mul/run_vck5000.lit index c7349d8844..4dcd76ba7c 100644 --- a/programming_examples/basic/vector_scalar_mul/run_vck5000.lit +++ b/programming_examples/basic/vector_scalar_mul/run_vck5000.lit @@ -1,10 +1,9 @@ -// (c) Copyright 2023 Advanced Micro Devices, Inc. +// (c) Copyright 2024 Advanced Micro Devices, Inc. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // -// REQUIRES: hsa +// REQUIRES: hsa, chess // -// RUN: xchesscc_wrapper aie -I %aietools/include -c %S/../../../aie_kernels/generic/scale.cc -o ./scale.o +// RUN: xchesscc_wrapper aie -I %aietools/include -c %S/scale.cc -o ./scale.o // RUN: %python %S/aie2.py xcvc1902 6 > ./aie.mlir // RUN: %python aiecc.py %VitisSysrootFlag% --host-target=%aieHostTargetTriplet% %link_against_hsa% ./aie.mlir -I%host_runtime_lib%/test_lib/include %extraAieCcFlags% %S/test_vck5000.cpp -o test.elf -L%host_runtime_lib%/test_lib/lib -ltest_lib -// RUN: %run_on_vck5000 ./test.elf - +// RU: %run_on_vck5000 ./test.elf