From 1c58bfd7fdd97ac2d6cf4cfbacb96e96bd4e1d12 Mon Sep 17 00:00:00 2001 From: James Edwards Date: Wed, 25 May 2016 10:39:06 -0500 Subject: [PATCH] Add hsa-rocr-dev packaging CMakeList.txt file. Change-Id: I1f6a0d4ad44aa7f20f43d43942719f668b620c36 [ROCm/ROCR-Runtime commit: ec6478e6936450d988d71039ec5a19e574168da5] --- .../packages/hsa-rocr-dev/CMakeLists.txt | 186 +++++++ .../runtime/packages/hsa-rocr-dev/LICENSE.txt | 37 ++ .../runtime/packages/hsa-rocr-dev/README.md | 78 +++ .../runtime/packages/hsa-rocr-dev/copyright | 38 ++ .../runtime/packages/hsa-rocr-dev/description | 1 + .../runtime/packages/hsa-rocr-dev/postinst | 19 + .../runtime/packages/hsa-rocr-dev/prerm | 19 + .../runtime/packages/hsa-rocr-dev/rpm_post | 1 + .../runtime/packages/hsa-rocr-dev/rpm_postun | 1 + .../packages/hsa-rocr-dev/sample/Makefile | 60 +++ .../hsa-rocr-dev/sample/vector_copy.c | 458 ++++++++++++++++++ .../hsa-rocr-dev/sample/vector_copy_base.brig | Bin 0 -> 3456 bytes .../sample/vector_copy_base.hsail | 63 +++ .../hsa-rocr-dev/sample/vector_copy_full.brig | Bin 0 -> 3456 bytes .../sample/vector_copy_full.hsail | 63 +++ 15 files changed, 1024 insertions(+) create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/CMakeLists.txt create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/LICENSE.txt create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/README.md create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/copyright create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/description create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/postinst create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/prerm create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_post create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_postun create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/Makefile create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy.c create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.brig create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.hsail create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_full.brig create mode 100644 projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_full.hsail diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/CMakeLists.txt b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/CMakeLists.txt new file mode 100644 index 0000000000..0bd41c9803 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/CMakeLists.txt @@ -0,0 +1,186 @@ +################################################################################ +## +## The University of Illinois/NCSA +## Open Source License (NCSA) +## +## Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +## +## Developed by: +## +## AMD Research and AMD HSA Software Development +## +## Advanced Micro Devices, Inc. +## +## www.amd.com +## +## Permission is hereby granted, free of charge, to any person obtaining a copy +## of this software and associated documentation files (the "Software"), to +## deal with the Software without restriction, including without limitation +## the rights to use, copy, modify, merge, publish, distribute, sublicense, +## and#or sell copies of the Software, and to permit persons to whom the +## Software is furnished to do so, subject to the following conditions: +## +## - Redistributions of source code must retain the above copyright notice, +## this list of conditions and the following disclaimers. +## - Redistributions in binary form must reproduce the above copyright +## notice, this list of conditions and the following disclaimers in +## the documentation and#or other materials provided with the distribution. +## - Neither the names of Advanced Micro Devices, Inc, +## nor the names of its contributors may be used to endorse or promote +## products derived from this Software without specific prior written +## permission. +## +## THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +## IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +## FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +## THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +## OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +## ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +## DEALINGS WITH THE SOFTWARE. +## +################################################################################ + +cmake_minimum_required(VERSION 2.8) + +set ( PROJECT_STRING hsa-rocr-dev ) + +project( ${PROJECT_STRING} ) + +list(APPEND CMAKE_MODULE_PATH "${HSA_OPENSOURCE_ROOT}/cmake_modules") + +include ( utils ) + +if(NOT DEFINED VERSION_STRING) + set (VERSION_STRING "1.0.0") +endif() + +parse_version(${VERSION_STRING}) + +set ( EXTERNAL_DIR ${CMAKE_CURRENT_BINARY_DIR} ) + +set ( RUNTIME_NAME "libhsa-runtime64" ) + +set ( RUNTIME_LIBRARY_SOURCE "${OUT_DIR}/lib/${RUNTIME_NAME}.so.${VERSION_STRING}" ) +set ( RUNTIME_LIBRARY_TARGET "${EXTERNAL_DIR}/hsa/lib/${RUNTIME_NAME}.so.${VERSION_STRING}" ) +set ( RUNTIME_LIBRARY_SONAME "${EXTERNAL_DIR}/hsa/lib/${RUNTIME_NAME}.so.1" ) +set ( RUNTIME_LIBRARY_LINKNAME "${EXTERNAL_DIR}/lib/${RUNTIME_NAME}.so" ) + +set ( PACKAGE_DIRECTORIES + "include" + "hsa/include" + "hsa/lib") + +set ( RUNTIME_HEADER_NAMES + "hsa.h" + "amd_hsa_common.h" + "amd_hsa_elf.h" + "amd_hsa_kernel_code.h" + "amd_hsa_queue.h" + "amd_hsa_signal.h" + "Brig.h" + "hsa_api_trace.h" + "hsa_ext_amd.h" + "hsa_ext_finalize.h" + "hsa_ext_image.h" + "hsa_ven_amd_loaded_code_object.h" + ) + +set ( RUNTIME_TOOLS_HEADER_NAMES + "hsa_ext_profiler.h" + "hsa_ext_debugger.h" + "amd_hsa_tools_interfaces.h" + ) + +set ( RUNTIME_TOOLS_INCLUDE_DIR "${AMD_CLOSED_SOURCE_ROOT}/drivers/hsa/runtime/tools/inc" ) + +set ( HSA_INCLUDE_LINK "${EXTERNAL_DIR}/hsa/include/hsa" ) + +set ( TOPLEVEL_INCLUDE_LINK "${EXTERNAL_DIR}/include/hsa" ) + +add_custom_command(OUTPUT ${PACKAGE_DIRECTORIES} WORKING_DIRECTORY ${EXTERNAL_DIR} + COMMAND mkdir -p include + COMMAND mkdir -p lib + COMMAND mkdir -p hsa/lib + COMMAND mkdir -p hsa/include ) + +add_custom_target (create_dirs DEPENDS ${PACKAGE_DIRECTORIES} ) + +add_custom_command( OUTPUT ${HSA_INCLUDE_LINK} WORKING_DIRECTORY ${EXTERNAL_DIR}/hsa/include + COMMAND ${CMAKE_COMMAND} -E create_symlink ../include hsa ) + +add_custom_command( OUTPUT ${TOPLEVEL_INCLUDE_LINK} WORKING_DIRECTORY ${EXTERNAL_DIR}/include + COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/include hsa ) + +add_custom_target (link_dirs DEPENDS create_dirs ${TOPLEVEL_INCLUDE_LINK} ${HSA_INCLUDE_LINK} ) + +add_custom_command( OUTPUT ${RUNTIME_LIBRARY_TARGET} WORKING_DIRECTORY ${EXTERNAL_DIR} + COMMAND ${CMAKE_COMMAND} -E copy ${RUNTIME_LIBRARY_SOURCE} ${RUNTIME_LIBRARY_TARGET} ) + +add_custom_command( OUTPUT ${RUNTIME_LIBRARY_SONAME} WORKING_DIRECTORY ${EXTERNAL_DIR}/hsa/lib + COMMAND ${CMAKE_COMMAND} -E create_symlink ${RUNTIME_NAME}.so.${VERSION_STRING} ${RUNTIME_NAME}.so.1 ) + +add_custom_command( OUTPUT ${RUNTIME_LIBRARY_LINKNAME} WORKING_DIRECTORY ${EXTERNAL_DIR}/lib + COMMAND ${CMAKE_COMMAND} -E create_symlink ../hsa/lib/${RUNTIME_NAME}.so.1 ${RUNTIME_NAME}.so ) + +foreach ( HEADER_FILE ${RUNTIME_HEADER_NAMES} ) + + set ( HEADER_TARGET "${EXTERNAL_DIR}/hsa/include/${HEADER_FILE}" ) + + add_custom_command( OUTPUT ${HEADER_TARGET} WORKING_DIRECTORY ${EXTERNAL_DIR} + COMMAND ${CMAKE_COMMAND} -E copy ${OUT_DIR}/include/${HEADER_FILE} ${HEADER_TARGET} ) + + list ( APPEND RUNTIME_HEADERS ${HEADER_TARGET} ) + +endforeach() + +foreach ( HEADER_FILE ${RUNTIME_TOOLS_HEADER_NAMES} ) + + set ( HEADER_TARGET "${EXTERNAL_DIR}/hsa/include/${HEADER_FILE}" ) + + add_custom_command( OUTPUT ${HEADER_TARGET} WORKING_DIRECTORY ${EXTERNAL_DIR} + COMMAND ${CMAKE_COMMAND} -E copy ${RUNTIME_TOOLS_INCLUDE_DIR}/${HEADER_FILE} ${HEADER_TARGET} ) + + list ( APPEND RUNTIME_HEADERS ${HEADER_TARGET} ) + +endforeach() + +add_custom_target (copy_headers DEPENDS ${RUNTIME_HEADERS} create_dirs) + +add_custom_target (copy_targets DEPENDS create_dirs + ${RUNTIME_LIBRARY_TARGET} + ${RUNTIME_LIBRARY_SONAME} + ${RUNTIME_LIBRARY_LINKNAME} + link_dirs + copy_headers + ) + +## Install directives +install (DIRECTORY ${EXTERNAL_DIR}/include/ DESTINATION include USE_SOURCE_PERMISSIONS ) +install (DIRECTORY ${EXTERNAL_DIR}/lib/ DESTINATION lib USE_SOURCE_PERMISSIONS ) +install (DIRECTORY ${EXTERNAL_DIR}/hsa/ DESTINATION hsa USE_SOURCE_PERMISSIONS ) +install (DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/sample DESTINATION hsa USE_SOURCE_PERMISSIONS ) +install (FILES ${CMAKE_CURRENT_SOURCE_DIR}/LICENSE.txt DESTINATION hsa ) +install (FILES ${CMAKE_CURRENT_SOURCE_DIR}/README.md DESTINATION hsa ) + +## Packaging directives +set ( CPACK_PACKAGE_NAME ${PROJECT_NAME} ) +set ( CPACK_PACKAGE_VENDOR "AMD" ) +set ( CPACK_PACKAGE_VERSION_MAJOR ${VERSION_MAJOR} ) +set ( CPACK_PACKAGE_VERSION_MINOR ${VERSION_MINOR} ) +set ( CPACK_PACKAGE_VERSION_PATCH ${VERSION_PATCH} ) +set ( CPACK_PACKAGE_CONTACT "James Edwards (JamesAdrian.Edwards@amd.com)" ) +set ( CPACK_PACKAGE_DESCRIPTION_SUMMARY "AMD Heterogeneous System Architecture HSA - Linux HSA Runtime for Boltzmann (ROCm) platforms" ) +set ( CPACK_PACKAGE_DESCRIPTION_FILE "${CMAKE_CURRENT_SOURCE_DIR}/description" ) +set ( CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/copyright" ) + +# Debian package specific variables +set ( CPACK_DEBIAN_PACKAGE_DEPENDS "hsa-ext-rocr-dev (=${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}), hsakmt-roct-dev (=2.0.0)" ) +set ( CPACK_DEBIAN_PACKAGE_HOMEPAGE "https://github.com/RadeonOpenCompute/ROCR-Runtime" ) +set ( CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${CMAKE_CURRENT_SOURCE_DIR}/postinst;${CMAKE_CURRENT_SOURCE_DIR}/prerm" ) + +# RPM package specific variables +set ( CPACK_RPM_PACKAGE_DEPENDS "hsa-ext-rocr-dev (=${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}), hsakmt-roct-dev (=2.0.0)" ) +set ( CPACK_RPM_PRE_INSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/rpm_post" ) +set ( CPACK_RPM_POST_UNINSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/rpm_postun" ) + +include ( CPack ) diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/LICENSE.txt b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/LICENSE.txt new file mode 100644 index 0000000000..8da0f93e08 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/LICENSE.txt @@ -0,0 +1,37 @@ +The University of Illinois/NCSA +Open Source License (NCSA) + +Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. + +Developed by: + + AMD Research and AMD HSA Software Development + + Advanced Micro Devices, Inc. + + www.amd.com + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to +deal with the Software without restriction, including without limitation +the rights to use, copy, modify, merge, publish, distribute, sublicense, +and/or sell copies of the Software, and to permit persons to whom the +Software is furnished to do so, subject to the following conditions: + + - Redistributions of source code must retain the above copyright notice, + this list of conditions and the following disclaimers. + - Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimers in + the documentation and/or other materials provided with the distribution. + - Neither the names of Advanced Micro Devices, Inc, + nor the names of its contributors may be used to endorse or promote + products derived from this Software without specific prior written + permission. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +DEALINGS WITH THE SOFTWARE. diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/README.md b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/README.md new file mode 100644 index 0000000000..b73327a533 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/README.md @@ -0,0 +1,78 @@ +### HSA Runtime API and runtime for Boltzmann + +This repository includes the user-mode API interfaces and libraries necessary for host applications to launch compute kernels to available HSA Boltzmann kernel agents. Reference source code for the core runtime is also available. + +Only the AMD/ATI Fiji(c) family of discrete GPUs are currently supported. + +#### Initial Target Platform Requirements + +* CPU: Intel(c) Haswell or newer, Core i5, Core i7, Xeon E3 v4 & v5; Xeon E5 v3 +* GPU: Fiji ASIC (AMD R9 Nano, R9 Fury and R9 Fury X) + +#### Source code + +The HSA core runtime source code for Boltzmann is located in the src subdirectory. Please consult the associated README.md file for contents and build instructions. + +#### Binaries for Ubuntu & Fedora and Installation Instructions + +Pre-built binaries are available for installation from the ROCm package repository. For ROCR, they include: + +Core runtime package: + +* HSA include files to support application development on the HSA runtime for Boltzmann +* A 64-bit version of AMD's HSA core runtime for Boltzmann + +Runtime extension package: + +* A 64-bit version of AMD's finalizer extension for Boltzmann +* A 64-bit version of AMD's runtime tools library + +The contents of these packages are installed in /opt/rocm/hsa and /opt/rocm by default. +The core runtime package depends on the hsakmt-roct-dev package + +Installation instructions can be found in the ROCm manifest repository README.md: + +https://github.com/RadeonOpenCompute/ROCm + +#### Infrastructure + +The HSA runtime is a thin, user-mode API that exposes the necessary interfaces to access and interact with graphics hardware driven by the AMDGPU driver set and the Boltzmann HSA kernel driver. Together they enable programmers to directly harness the power of AMD discrete graphics devices by allowing host applications to launch compute kernels directly to the graphics hardware. + +The capabilities expressed by the HSA Runtime API are: + +* Error handling +* Runtime initialization and shutdown +* System and agent information +* Signals and synchronization +* Architected dispatch +* Memory management +* HSA runtime fits into a typical software architecture stack. + +The HSA runtime provides direct access to the graphics hardware to give the programmer more control of the execution. Some examples of low level hardware access is the support of one or more user mode queues provides programmers with a low-latency kernel dispatch interface, allowing them to develop customized dispatch algorithms specific to their application. + +The HSA Architected Queuing Language is an open standard, defined by the HSA Foundation, specifying the packet syntax used to control supported AMD/ATI Radeon (c) graphics devices. The AQL language supports several packet types, including packets that can command the hardware to automatically resolve inter-packet dependencies (barrier AND & barrier OR packet), kernel dispatch packets and agent dispatch packets. + +In addition to user mode queues and AQL, the HSA runtime exposes various virtual address ranges that can be accessed by one or more of the system's graphics devices, and possibly the host. The exposed virtual address ranges either support a fine grained or a coarse grained access. Updates to memory in a fine grained region are immediately visible to all devices that can access it, but only one device can have access to a coarse grained allocation at a time. Ownership of a coarse grained region can be changed using the HSA runtime memory APIs, but this transfer of ownership must be explicitly done by the host application. + +Programmers should consult the HSA Runtime Programmer's Reference Manual for a full description of the HSA Runtime APIs, AQL and the HSA memory policy. + +#### Sample + +The simplest way to check if the kernel, runtime and base development environment are installed correctly is to run a simple sample. A modified version of the vector_copy sample was taken from the HSA-Runtime-AMD repository and added to the ROCR repository to facilitate this. Build the sample and run it, using this series of commands: + +cd ROCR-Runtime/sample && make && ./vector_copy + +If the sample runs without generating errors, the installation is complete. + +#### Known Issues + +* The image extension is currently not supported for discrete GPUs. An image extension library is not provided in the binary package. The standard hsa_ext_image.h extension include file is provided for reference. +* Each HSA process creates and internal DMA queue, but there is a system-wide limit of four DMA queues. The fifths simultaneous HSA process will fail hsa_init() with HSA_STATUS_ERROR_OUT_OF_RESOURCES. To run an unlimited number of simultaneous HSA processes, set the environment variable HSA_ENABLE_SDMA=0. + +#### Disclaimer + +The information contained herein is for informational purposes only, and is subject to change without notice. While every precaution has been taken in the preparation of this document, it may contain technical inaccuracies, omissions and typographical errors, and AMD is under no obligation to update or otherwise correct this information. Advanced Micro Devices, Inc. makes no representations or warranties with respect to the accuracy or completeness of the contents of this document, and assumes no liability of any kind, including the implied warranties of noninfringement, merchantability or fitness for particular purposes, with respect to the operation or use of AMD hardware, software or other products described herein. No license, including implied or arising by estoppel, to any intellectual property rights is granted by this document. Terms and limitations applicable to the purchase or use of AMD's products are as set forth in a signed agreement between the parties or in AMD's Standard Terms and Conditions of Sale. + +AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. + +Copyright (c) 2014-2016 Advanced Micro Devices, Inc. All rights reserved. diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/copyright b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/copyright new file mode 100644 index 0000000000..1fc00dc10a --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/copyright @@ -0,0 +1,38 @@ +The University of Illinois/NCSA +Open Source License (NCSA) + +Copyright (c) 2014-2016, Advanced Micro Devices, Inc. All rights reserved. + +Developed by: + + AMD Research and AMD HSA Software Development + + Advanced Micro Devices, Inc. + + www.amd.com + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to +deal with the Software without restriction, including without limitation +the rights to use, copy, modify, merge, publish, distribute, sublicense, +and/or sell copies of the Software, and to permit persons to whom the +Software is furnished to do so, subject to the following conditions: + + - Redistributions of source code must retain the above copyright notice, + this list of conditions and the following disclaimers. + - Redistributions in binary form must reproduce the above copyright + notice, this list of conditions and the following disclaimers in + the documentation and/or other materials provided with the distribution. + - Neither the names of Advanced Micro Devices, Inc, + nor the names of its contributors may be used to endorse or promote + products derived from this Software without specific prior written + permission. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +DEALINGS WITH THE SOFTWARE. + diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/description b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/description new file mode 100644 index 0000000000..553d517053 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/description @@ -0,0 +1 @@ +This package includes the user-mode runtime necessary for host applications to launch compute kernels to available HSA and Boltzmann components. This version is consistent with the 1.0 Final HSA Runtime Programmer's Reference Manual and targets AMD AMD Fiji ASICS on supported platforms. diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/postinst b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/postinst new file mode 100644 index 0000000000..769a72e462 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/postinst @@ -0,0 +1,19 @@ +#/bin/bash + +set -e + +do_ldconfig() { + echo /opt/rocm/hsa/lib > /etc/ld.so.conf.d/hsa-rocr-dev.conf && ldconfig +} + +case "$1" in + configure) + do_ldconfig + ;; + abort-upgrade|abort-remove|abort-deconfigure) + echo "$1" + ;; + *) + exit 0 + ;; +esac diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/prerm b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/prerm new file mode 100644 index 0000000000..2b7d50a825 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/prerm @@ -0,0 +1,19 @@ +#!/bin/bash + +set -e + +rm_ldconfig() { + rm -f /etc/ld.so.conf.d/hsa-rocr-dev.conf && ldconfig +} + +case "$1" in + remove) + rm_ldconfig + ;; + purge) + ;; + *) + exit 0 + ;; +esac + diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_post b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_post new file mode 100644 index 0000000000..5a9aaac79f --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_post @@ -0,0 +1 @@ +echo /opt/rocm/hsa/lib > /etc/ld.so.conf.d/hsa-rocr-dev.conf && ldconfig diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_postun b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_postun new file mode 100644 index 0000000000..a925febfc4 --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/rpm_postun @@ -0,0 +1 @@ +rm -f /etc/ld.so.conf.d/hsa-rocr-dev.conf && ldconfig diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/Makefile b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/Makefile new file mode 100644 index 0000000000..218dddaacf --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/Makefile @@ -0,0 +1,60 @@ +################################################################################ +## +## The University of Illinois/NCSA +## Open Source License (NCSA) +## +## Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +## +## Developed by: +## +## AMD Research and AMD HSA Software Development +## +## Advanced Micro Devices, Inc. +## +## www.amd.com +## +## Permission is hereby granted, free of charge, to any person obtaining a copy +## of this software and associated documentation files (the "Software"), to +## deal with the Software without restriction, including without limitation +## the rights to use, copy, modify, merge, publish, distribute, sublicense, +## and#or sell copies of the Software, and to permit persons to whom the +## Software is furnished to do so, subject to the following conditions: +## +## - Redistributions of source code must retain the above copyright notice, +## this list of conditions and the following disclaimers. +## - Redistributions in binary form must reproduce the above copyright +## notice, this list of conditions and the following disclaimers in +## the documentation and#or other materials provided with the distribution. +## - Neither the names of Advanced Micro Devices, Inc, +## nor the names of its contributors may be used to endorse or promote +## products derived from this Software without specific prior written +## permission. +## +## THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +## IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +## FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +## THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +## OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +## ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +## DEALINGS WITH THE SOFTWARE. +## +################################################################################ + +LFLAGS= -Wl,--unresolved-symbols=ignore-in-shared-libs + +CC := gcc + +C_FILES := $(wildcard *.c) + +OBJ_FILES := $(notdir $(C_FILES:.c=.o)) + +all: vector_copy + +vector_copy: $(OBJ_FILES) + $(CC) $(LFLAGS) $(OBJ_FILES) -L/opt/rocm/lib -lhsa-runtime64 -o vector_copy + +%.o: %.c + $(CC) -c -I/opt/rocm/include -o $@ $< -std=c99 + +clean: + rm -rf *.o vector_copy diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy.c b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy.c new file mode 100644 index 0000000000..3179de415b --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy.c @@ -0,0 +1,458 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +#include +#include +#include +#include +#include "hsa/hsa.h" +#include "hsa/hsa_ext_finalize.h" + +#define check(msg, status) \ +if (status != HSA_STATUS_SUCCESS) { \ + printf("%s failed.\n", #msg); \ + exit(1); \ +} else { \ + printf("%s succeeded.\n", #msg); \ +} + +/* + * Loads a BRIG module from a specified file. This + * function does not validate the module. + */ +int load_module_from_file(const char* file_name, hsa_ext_module_t* module) { + int rc = -1; + + FILE *fp = fopen(file_name, "rb"); + + rc = fseek(fp, 0, SEEK_END); + + size_t file_size = (size_t) (ftell(fp) * sizeof(char)); + + rc = fseek(fp, 0, SEEK_SET); + + char* buf = (char*) malloc(file_size); + + memset(buf,0,file_size); + + size_t read_size = fread(buf,sizeof(char),file_size,fp); + + if(read_size != file_size) { + free(buf); + } else { + rc = 0; + *module = (hsa_ext_module_t) buf; + } + + fclose(fp); + + return rc; +} + +/* + * Determines if the given agent is of type HSA_DEVICE_TYPE_GPU + * and sets the value of data to the agent handle if it is. + */ +static hsa_status_t get_gpu_agent(hsa_agent_t agent, void *data) { + hsa_status_t status; + hsa_device_type_t device_type; + status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type); + if (HSA_STATUS_SUCCESS == status && HSA_DEVICE_TYPE_GPU == device_type) { + hsa_agent_t* ret = (hsa_agent_t*)data; + *ret = agent; + return HSA_STATUS_INFO_BREAK; + } + return HSA_STATUS_SUCCESS; +} + +/* + * Determines if a memory region can be used for kernarg + * allocations. + */ +static hsa_status_t get_kernarg_memory_region(hsa_region_t region, void* data) { + hsa_region_segment_t segment; + hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment); + if (HSA_REGION_SEGMENT_GLOBAL != segment) { + return HSA_STATUS_SUCCESS; + } + + hsa_region_global_flag_t flags; + hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags); + if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) { + hsa_region_t* ret = (hsa_region_t*) data; + *ret = region; + return HSA_STATUS_INFO_BREAK; + } + + return HSA_STATUS_SUCCESS; +} + +/* + * Determines if a memory region can be used for fine grained + * allocations. + */ +static hsa_status_t get_fine_grained_memory_region(hsa_region_t region, void* data) { + hsa_region_segment_t segment; + hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment); + if (HSA_REGION_SEGMENT_GLOBAL != segment) { + return HSA_STATUS_SUCCESS; + } + + hsa_region_global_flag_t flags; + hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags); + if (flags & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED) { + hsa_region_t* ret = (hsa_region_t*) data; + *ret = region; + return HSA_STATUS_INFO_BREAK; + } + + return HSA_STATUS_SUCCESS; +} + +int main(int argc, char **argv) { + hsa_status_t err; + + err = hsa_init(); + check(Initializing the hsa runtime, err); + + /* + * Determine if the finalizer 1.0 extension is supported. + */ + bool support; + + err = hsa_system_extension_supported(HSA_EXTENSION_FINALIZER, 1, 0, &support); + + check(Checking finalizer 1.0 extension support, err); + + /* + * Generate the finalizer function table. + */ + hsa_ext_finalizer_1_00_pfn_t table_1_00; + + err = hsa_system_get_extension_table(HSA_EXTENSION_FINALIZER, 1, 0, &table_1_00); + + check(Generating function table for finalizer, err); + + /* + * Iterate over the agents and pick the gpu agent using + * the get_gpu_agent callback. + */ + hsa_agent_t agent; + err = hsa_iterate_agents(get_gpu_agent, &agent); + if(err == HSA_STATUS_INFO_BREAK) { + err = HSA_STATUS_SUCCESS; + } else { + /* + * No GPU agent was found. + */ + err = HSA_STATUS_ERROR; + } + check(Getting a gpu agent, err); + + /* + * Query the name of the agent. + */ + char name[64] = { 0 }; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name); + check(Querying the agent name, err); + printf("The agent name is %s.\n", name); + + /* + * Query the maximum size of the queue. + */ + uint32_t queue_size = 0; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size); + check(Querying the agent maximum queue size, err); + printf("The maximum queue size is %u.\n", (unsigned int) queue_size); + + /* + * Create a queue using the maximum size. + */ + hsa_queue_t* queue; + err = hsa_queue_create(agent, queue_size, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &queue); + check(Creating the queue, err); + + /* + * Obtain the agent's machine model + */ + hsa_machine_model_t machine_model; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model); + check("Obtaining machine model",err); + + /* + * Obtain the agent's profile + */ + hsa_profile_t profile; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &profile); + check("Getting agent profile",err); + + /* + * Load the BRIG binary. + */ + hsa_ext_module_t module; + if(HSA_PROFILE_FULL == profile) { + load_module_from_file("vector_copy_full.brig",&module); + } else { + load_module_from_file("vector_copy_base.brig",&module); + } + + /* + * Create hsa program. + */ + hsa_ext_program_t program; + memset(&program,0,sizeof(hsa_ext_program_t)); + err = table_1_00.hsa_ext_program_create(machine_model, profile, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, NULL, &program); + check(Create the program, err); + + /* + * Add the BRIG module to hsa program. + */ + err = table_1_00.hsa_ext_program_add_module(program, module); + check(Adding the brig module to the program, err); + + /* + * Determine the agents ISA. + */ + hsa_isa_t isa; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_ISA, &isa); + check(Query the agents isa, err); + + /* + * Finalize the program and extract the code object. + */ + hsa_ext_control_directives_t control_directives; + memset(&control_directives, 0, sizeof(hsa_ext_control_directives_t)); + hsa_code_object_t code_object; + err = table_1_00.hsa_ext_program_finalize(program, isa, 0, control_directives, "", HSA_CODE_OBJECT_TYPE_PROGRAM, &code_object); + check(Finalizing the program, err); + + /* + * Destroy the program, it is no longer needed. + */ + err=table_1_00.hsa_ext_program_destroy(program); + check(Destroying the program, err); + + /* + * Create the empty executable. + */ + hsa_executable_t executable; + err = hsa_executable_create(profile, HSA_EXECUTABLE_STATE_UNFROZEN, "", &executable); + check(Create the executable, err); + + /* + * Load the code object. + */ + err = hsa_executable_load_code_object(executable, agent, code_object, ""); + check(Loading the code object, err); + + /* + * Freeze the executable; it can now be queried for symbols. + */ + err = hsa_executable_freeze(executable, ""); + check(Freeze the executable, err); + + /* + * Extract the symbol from the executable. + */ + hsa_executable_symbol_t symbol; + err = hsa_executable_get_symbol(executable, NULL, "&__vector_copy_kernel", agent, 0, &symbol); + check(Extract the symbol from the executable, err); + + /* + * Extract dispatch information from the symbol + */ + uint64_t kernel_object; + uint32_t kernarg_segment_size; + uint32_t group_segment_size; + uint32_t private_segment_size; + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object); + check(Extracting the symbol from the executable, err); + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &kernarg_segment_size); + check(Extracting the kernarg segment size from the executable, err); + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_segment_size); + check(Extracting the group segment size from the executable, err); + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &private_segment_size); + check(Extracting the private segment from the executable, err); + + /* + * Create a signal to wait for the dispatch to finish. + */ + hsa_signal_t signal; + err=hsa_signal_create(1, 0, NULL, &signal); + check(Creating a HSA signal, err); + + hsa_region_t finegrained_region; + finegrained_region.handle=(uint64_t)-1; + hsa_agent_iterate_regions(agent, get_fine_grained_memory_region, &finegrained_region); + err = (finegrained_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS; + check(Finding a fine grained memory region, err); + + /* + * Allocate and initialize the kernel arguments from the fine + * grained memory region. + */ + char* in; + err=hsa_memory_allocate(finegrained_region, 1024*1024*4, (void*) &in); + check(Allocating argument memory for input parameter, err); + memset(in, 1, 1024*1024*4); + + char* out; + err=hsa_memory_allocate(finegrained_region, 1024*1024*4, (void*) &out); + check(Allocating argument memory for output parameter, err); + memset(out, 0, 1024*1024*4); + + struct __attribute__ ((aligned(16))) args_t { + void* in; + void* out; + } args; + + args.in=in; + args.out=out; + + /* + * Find a memory region that supports kernel arguments. + */ + hsa_region_t kernarg_region; + kernarg_region.handle=(uint64_t)-1; + hsa_agent_iterate_regions(agent, get_kernarg_memory_region, &kernarg_region); + err = (kernarg_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS; + check(Finding a kernarg memory region, err); + void* kernarg_address = NULL; + + /* + * Allocate the kernel argument buffer from the correct region. + */ + err = hsa_memory_allocate(kernarg_region, kernarg_segment_size, &kernarg_address); + check(Allocating kernel argument memory buffer, err); + memcpy(kernarg_address, &args, sizeof(args)); + + /* + * Obtain the current queue write index. + */ + uint64_t index = hsa_queue_load_write_index_relaxed(queue); + + /* + * Write the aql packet at the calculated queue index address. + */ + const uint32_t queueMask = queue->size - 1; + hsa_kernel_dispatch_packet_t* dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(queue->base_address))[index&queueMask]); + + dispatch_packet->setup |= 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + dispatch_packet->workgroup_size_x = (uint16_t)256; + dispatch_packet->workgroup_size_y = (uint16_t)1; + dispatch_packet->workgroup_size_z = (uint16_t)1; + dispatch_packet->grid_size_x = (uint32_t) (1024*1024); + dispatch_packet->grid_size_y = 1; + dispatch_packet->grid_size_z = 1; + dispatch_packet->completion_signal = signal; + dispatch_packet->kernel_object = kernel_object; + dispatch_packet->kernarg_address = (void*) kernarg_address; + dispatch_packet->private_segment_size = private_segment_size; + dispatch_packet->group_segment_size = group_segment_size; + + uint16_t header = 0; + header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; + header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + header |= HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + + __atomic_store_n((uint16_t*)(&dispatch_packet->header), header, __ATOMIC_RELEASE); + + /* + * Increment the write index and ring the doorbell to dispatch the kernel. + */ + hsa_queue_store_write_index_relaxed(queue, index+1); + hsa_signal_store_relaxed(queue->doorbell_signal, index); + check(Dispatching the kernel, err); + + /* + * Wait on the dispatch completion signal until the kernel is finished. + */ + hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + /* + * Validate the data in the output buffer. + */ + int valid=1; + int fail_index=0; + for(int i=0; i<1024*1024; i++) { + if(out[i]!=in[i]) { + fail_index=i; + valid=0; + break; + } + } + + if(valid) { + printf("Passed validation.\n"); + } else { + printf("VALIDATION FAILED!\nBad index: %d\n", fail_index); + } + + /* + * Cleanup all allocated resources. + */ + err = hsa_memory_free(kernarg_address); + check(Freeing kernel argument memory buffer, err); + + err=hsa_signal_destroy(signal); + check(Destroying the signal, err); + + err=hsa_executable_destroy(executable); + check(Destroying the executable, err); + + err=hsa_code_object_destroy(code_object); + check(Destroying the code object, err); + + err=hsa_queue_destroy(queue); + check(Destroying the queue, err); + + err = hsa_memory_free(in); + check(Freeing in argument memory buffer, err); + + err = hsa_memory_free(out); + check(Freeing out argument memory buffer, err); + + err=hsa_shut_down(); + check(Shutting down the runtime, err); + + return 0; +} diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.brig b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.brig new file mode 100644 index 0000000000000000000000000000000000000000..563960cd63aceaa08b09af3e364a7b0324af31ed GIT binary patch literal 3456 zcmb_ey>BB$6rbb*7&wrTLIVeS3IW|oOA-(WbfUOTyp9)aueIKAchX6@_3k7duy>=~ zb)rZVS5Qz=UP(bgNpS@Q1qB5K1qCI40tMw26#U-Itj{qA1!9!Rn|U9<_c3qA+ZnpF z8~Xbz!i;;**~R}rE&*d~aqe0A!_#JQ_p+fk?4Gb~V2;c9FpaC&NQD?+vHd@fRgEdc zhNcvaB^_kyR2HSG&M7}4e=<>7uFCEH-q01|6{8oN$c%>htVkqvRU)%e(pA>D26@lu z?Byrtg*qNrbTzp~Z|}VI&YSr6u0w8m8fS@2si%@6r$3>h~GIi*vkdkYhjN{^1 zI#lI|_nb~peHi4UDpr}wjwz-DwGFS&76re21oas*adSo$bkNG*RW)p;WPl(X4`I$xPi(j_8SEnvy@JR9dz&THRb)fM2=IGPn1{G(#fOLNpyyZ=Nc6>~U^qkJ;S z&p6#A&r-!TD6fmI!J;?OX7kut#Br78QY3jQX*w&pete7?E*|BlQWwMAi!85@tMh@y zBPby*6s6K#F?VUwyATijCX-m9l;xJq?J@Dd??`2FagI!iX&ui=k*70UT+VPQB+;J7 zVuSl{F`-d}=+Cj()5W&5hnhoD&=h@JsCZHmLOWF{j%41Civ`9*@ZP-keTfEzF=J=( z)O5waJS~p!=ka9drwewfD*1>_8|Fockj`l&xqEec$SegHm$l#!r+FoTW&0D)kEtjy zQsjXb3pZQ09ZqxJarK2MPh_HwR01OfxU)i4mCRn4$EJH&KNn> z$RE-m3_kW-UW+!|A@&;%-SneQa1aqRLbo5?rl3u3|2BQ(_gfBmw+5j%98wSh)4n5q zZ_xEU2>t$E_n_tX_h}aZ{UD;Q-}57&M**>Ci|Km+Xj9J%_d0Ol?)qIny6uP$P3&zy z>a&k_5Ryv+H;nwfgRUFW;2<0XLl5D!fZq4}?GVnqp4X4I;2b>iKE{TI9k+|iW$zbz z-V6Fs=p(Ek@Kaf2asoboZ@i96xW*LE_{*R-9;bVVUy9;fyaoC8Htp`dwsU>y_w;XD=)G>N zgvN)d{$AMxb?o|m&>NWSpJTs_$>+e|Ft;!pm?vnn>d8u1 zM5DQ*Wvk)=_P?_ndBY5{=bK$?wuSve@Sd&y}Dvb zz`PE8?--s=fH9A;`T)FotRr6@o0!jHa)07m7GHxYv-N0*XSlTx)7lP>Ts%Nb%tN^} zW!d7M`L*t^bKMohDRzFb{-C0Vv0JhBHLdYhwGE%o;<)H*^@8ukSlA!i{kmYxwvTj7 MOe}u5ptHFA3v;;3s{jB1 literal 0 HcmV?d00001 diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.hsail b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.hsail new file mode 100644 index 0000000000..7eeeb471db --- /dev/null +++ b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_base.hsail @@ -0,0 +1,63 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +module &m:1:0:$base:$large:$default; + +decl prog function &abort()(); + +prog kernel &__vector_copy_kernel( + kernarg_u64 %in, + kernarg_u64 %out) +{ +@__vector_copy_kernel_entry: + // BB#0: // %entry + workitemabsid_u32 $s0, 0; + cvt_s64_s32 $d0, $s0; + shl_u64 $d0, $d0, 2; + ld_kernarg_align(8)_width(all)_u64 $d1, [%out]; + add_u64 $d1, $d1, $d0; + ld_kernarg_align(8)_width(all)_u64 $d2, [%in]; + add_u64 $d0, $d2, $d0; + ld_global_u32 $s0, [$d0]; + st_global_u32 $s0, [$d1]; + ret; +}; diff --git a/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_full.brig b/projects/rocr-runtime/runtime/packages/hsa-rocr-dev/sample/vector_copy_full.brig new file mode 100644 index 0000000000000000000000000000000000000000..bd042dbffecdd9738d50f931102d1a987cab55ea GIT binary patch literal 3456 zcmb_ey>BB$6rbb*7&wrTLIVeS3IW|oOA-(WbfUOTyp9)aueIKAchX6@_3k7duy>=~ zb)rZVS5Qz=UP(bgNpS@Q1qB5K1qCI40tMw26#U-Itj{qA1!9!Rn|U9<_c3qA+ZnpF z8~Xbz!i;;**~R}rE&*d~aqe0A!_#JQ_p+fk?4Gb~V2;c9FpaC&NQD?+vHd@fRgEdc zhNcvaB^_kyR2HSG&M7}4e=<>7uFCEH-q01|6{8oN$c%>htVkqvRU)%e(pA>D26@lu z?Byrtg*qNrbTzp~Z|}VI&YSr6u0w8m8fS@2si%@6r$3>h~GIi*vkdkYhjN{^1 zI#lI|_nb~peHi4UDpr}wjwz-DwGFS&76re21oas*adSo$bkNG*RW)p;WPl(X4`I$xPi(j_8SEnvy@JR9dz&THRb)fM2=IGPn1{G(#fOLNpyyZ=Nc6>~U^qkJ;S z&p6#A&r-!TD6fmI!J;?OX7kut#Br78QY3jQX*w&pete7?E*|BlQWwMAi!85@tMh@y zBPby*6s6K#F?VUwyATijCX-m9l;xJq?J@Dd??`2FagI!iX&ui=k*70UT+VPQB+;J7 zVuSl{F`-d}=+Cj()5W&5hnhoD&=h@JsCZHmLOWF{j%41Civ`9*@ZP-keTfEzF=J=( z)O5waJS~p!=ka9drwewfD*1>_8|Fockj`l&xqEec$SegHm$l#!r+FoTW&0D)kEtjy zQsjXb3pZQ09ZqxJarK2MPh_HwR01OfxU)i4mCRn4$EJH&KNn> z$RE-m3_kW-UW+!|A@&;%-SneQa1aqRLbo5?rl3u3|2BQ(_gfBmw+5j%98wSh)4n5q zZ_xEU2>t$E_n_tX_h}aZ{UD;Q-}57&M**>Ci|Km+Xj9J%_d0Ol?)qIny6uP$P3&zy z>a&k_5Ryv+H;nwfgRUFW;2<0XLl5D!fZq4}?GVnqp4X4I;2b>iKE{TI9k+|iW$zbz z-V6Fs=p(Ek@Kaf2asoboZ@i96xW*LE_{*R-9;bVVUy9;fyaoC8Htp`dwsU>y_w;XD=)G>N zgvN)d{$AMxb?o|m&>NWSpJTs_$>+e|Ft;!pm?TR#FNcUE;`122M zG-cW1p82)zuXEiM#3^=uvHqZ