commit: 5d15eda65acc6e32daf37f651fa8f216065162f2
Author: Craig Andrews <candrews <AT> gentoo <DOT> org>
AuthorDate: Fri Aug 9 21:47:55 2019 +0000
Commit: Craig Andrews <candrews <AT> gentoo <DOT> org>
CommitDate: Wed Aug 14 00:27:12 2019 +0000
URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=5d15eda6
dev-libs/rocm-opencl-runtime: Radeon Open Compute OpenCL Compatible Runtime
Package-Manager: Portage-2.3.71, Repoman-2.3.17
Signed-off-by: Craig Andrews <candrews <AT> gentoo.org>
dev-libs/rocm-opencl-runtime/Manifest | 2 +
...pencl-runtime-2.6.0-unbundle-dependencies.patch | 1233 ++++++++++++++++++++
dev-libs/rocm-opencl-runtime/metadata.xml | 14 +
.../rocm-opencl-runtime-2.6.0.ebuild | 50 +
4 files changed, 1299 insertions(+)
diff --git a/dev-libs/rocm-opencl-runtime/Manifest
b/dev-libs/rocm-opencl-runtime/Manifest
new file mode 100644
index 00000000000..7d10d714262
--- /dev/null
+++ b/dev-libs/rocm-opencl-runtime/Manifest
@@ -0,0 +1,2 @@
+DIST OpenCL-ICD-Loader-bc9728edf8cace79cf33bf75560be88fc2432dc4.tar.gz 62858
BLAKE2B
931fb7ade12debda99512d8b9fbfb987f47e3b68568863e480de406f842a014d39f0321c44420fb8b68b8757d49cc02cadf82e767c9d2b71e53bdb8e9763cace
SHA512
2ab9035fbc0a80a529d1b2f454b0b9a263b6eff601121750e81e2e32581f06189566cdcf201ab741ff3f1491ead825a7858000901f8042af67b1304be622e9fa
+DIST rocm-opencl-runtime-2.6.0.tar.gz 986155 BLAKE2B
43469ec3a368e2ee9716a38d08df48c28165a233bd80de54f286011df25782a197e252e8e41ffa8f59eecf93d6beae5434497e8d4b648676252d32a0859ba59a
SHA512
bbbbd29bf23f93440135786600cca28a806714bb207b7ad0947d26471fa38470820e7801acc84380000d43d55ff2ec1d95cee9e64b500628b6a3b026744c67df
diff --git
a/dev-libs/rocm-opencl-runtime/files/rocm-opencl-runtime-2.6.0-unbundle-dependencies.patch
b/dev-libs/rocm-opencl-runtime/files/rocm-opencl-runtime-2.6.0-unbundle-dependencies.patch
new file mode 100644
index 00000000000..379dd11e9f9
--- /dev/null
+++
b/dev-libs/rocm-opencl-runtime/files/rocm-opencl-runtime-2.6.0-unbundle-dependencies.patch
@@ -0,0 +1,1233 @@
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index ebdfc25..fb1c3eb 100644
+--- a/CMakeLists.txt
++++ b/CMakeLists.txt
+@@ -21,6 +21,8 @@ endif()
+
+ project(OpenCL-ROCm)
+
++include (GNUInstallDirs)
++
+ # Add path for custom modules
+ set(CMAKE_MODULE_PATH
+ ${CMAKE_MODULE_PATH}
+@@ -38,38 +40,28 @@ set(CLANG_ENABLE_STATIC_ANALYZER OFF CACHE BOOL "")
+ # override default option value in library and driver
+ set(GENERIC_IS_ZERO ON CACHE BOOL ON FORCE)
+
+-add_subdirectory(compiler/llvm EXCLUDE_FROM_ALL)
+-
+-find_package(LLVM REQUIRED CONFIG PATHS ${CMAKE_BINARY_DIR}/compiler/llvm
NO_DEFAULT_PATH)
++find_package(LLVM REQUIRED CONFIG PATHS ${LLVM_DIR} "/opt/rocm/llvm"
NO_DEFAULT_PATH)
+
+ list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
+ include(AddLLVM)
+
+ add_definitions(${LLVM_DEFINITIONS})
+-# TODO: add find_package for Clang and lld, and also use LLVM/Clang variables
got from their config
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/llvm/tools/clang/include)
+-include_directories(${CMAKE_BINARY_DIR}/compiler/llvm/tools/clang/include)
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/llvm/tools/lld/include)
+-
+-# TODO: move AMDGPU.h header to include folder
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/llvm/lib/Target/AMDGPU)
+-include_directories(${CMAKE_BINARY_DIR}/compiler/llvm/lib/Target/AMDGPU)
+-
+-if(${USE_COMGR_LIBRARY} MATCHES "yes")
+- set(COMGR_DYN_DLL "yes")
+- add_definitions(-DCOMGR_DYN_DLL="yes")
+- include_directories(${CMAKE_SOURCE_DIR}/compiler/llvm/include)
+- add_definitions(-DUSE_COMGR_LIBRARY)
+-else()
+- add_subdirectory(compiler/driver EXCLUDE_FROM_ALL)
+-endif()
++set(USE_COMGR_LIBRARY "yes")
++find_package(amd_comgr REQUIRED CONFIG)
++add_definitions(-DUSE_COMGR_LIBRARY)
++FOREACH(DIR ${LLVM_INCLUDE_DIRS})
++ include_directories("${DIR}")
++ include_directories("${DIR}/clang")
++ include_directories("${DIR}/lld")
++ # TODO: move AMDGPU.h header to include folder
++ include_directories("${DIR}/llvm/Target/AMDGPU")
++ENDFOREACH()
+
+ set(BUILD_HC_LIB OFF CACHE BOOL "")
+ set(ROCM_DEVICELIB_INCLUDE_TESTS OFF CACHE BOOL "")
+ set(AMDGCN_TARGETS_LIB_LIST "AMDGCN_LIB_TARGETS")
+ set(AMDGCN_TARGETS_LIB_DEPS "AMDGCN_DEP_TARGETS")
+ set(AMDGPU_TARGET_TRIPLE "amdgcn-amd-amdhsa-amdgizcl")
+-add_subdirectory(library/amdgcn EXCLUDE_FROM_ALL)
+
+ add_subdirectory(compiler/lib/loaders/elf/utils/libelf)
+
+@@ -82,28 +74,6 @@ set(OPENCL_INCLUDE_DIRS
+ ${CMAKE_SOURCE_DIR}/api/opencl/khronos/headers/opencl2.2)
+
+ add_subdirectory(api/opencl/amdocl)
+-add_subdirectory(api/opencl/khronos/icd)
+-
+-add_subdirectory(tools/clinfo)
+-
+-install(PROGRAMS $<TARGET_FILE:clang> $<TARGET_FILE:lld>
+- DESTINATION bin/x86_64
+- COMPONENT libraries)
+-
+-install(PROGRAMS $<TARGET_FILE:OpenCL>
+- DESTINATION lib/x86_64
+- COMPONENT applications)
+-
+-install(PROGRAMS $<TARGET_LINKER_FILE:OpenCL>
+- DESTINATION lib/x86_64
+- COMPONENT libraries)
+-
+-install(DIRECTORY
+- "${CMAKE_CURRENT_SOURCE_DIR}/api/opencl/khronos/headers/opencl2.2/CL"
+- DESTINATION include
+- COMPONENT libraries
+- USE_SOURCE_PERMISSIONS
+- PATTERN cl_egl.h EXCLUDE)
+
+ foreach(AMDGCN_LIB_TARGET ${AMDGCN_LIB_TARGETS})
+ get_target_property(lib_file_name ${AMDGCN_LIB_TARGET} ARCHIVE_OUTPUT_NAME)
+diff --git a/README.md b/README.md
+index a834965..e319be2 100644
+--- a/README.md
++++ b/README.md
+@@ -7,13 +7,6 @@ Developer preview Version 2 of the new
+ * Supports offline ahead of time compilation today; during the Beta phase we
will add in-process/in-memory compilation.
+
+
+-## GETTING REPO
+-
+-Repo is a git wrapper that manages a collection of git repositories. Install
this tool and add it to the command search PATH:
+-
+- curl https://storage.googleapis.com/git-repo-downloads/repo > ~/bin/repo
+- chmod a+x ~/bin/repo
+-
+ ## GETTING THE SOURCE CODE
+
+ Main OpenCL™ Compatible Components:
+@@ -26,11 +19,6 @@ Main OpenCL™ Compatible Components:
+ * https://github.com/RadeonOpenCompute/lld
+ * https://github.com/KhronosGroup/OpenCL-ICD-Loader
+
+-Download the git projects with the following commands:
+-
+- ~/bin/repo init -u
https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git -b master -m
opencl.xml
+- ~/bin/repo sync
+-
+ ## INSTALL ROCm
+
+ Follow the instructions at https://rocm.github.io/install.html to install
ROCm.
+@@ -45,6 +33,11 @@ Copy the amdocl64.icd file to /etc/OpenCL/vendors
+
+ To install additional dependencies:
+
++* ROCm-OpenCL-Runtime
++* ROCm-OpenCL-Driver
++* ROC versions of LLVM, Clang, and lld
++* ROCm-Device-Libs
++* OpenCL-ICD-Loader
+ * OCaml
+ * findlib
+ * A Python 2 environment or active virtualenv with the Microsoft Z3 package
+diff --git a/api/opencl/amdocl/CMakeLists.txt
b/api/opencl/amdocl/CMakeLists.txt
+index 91bd42f..8606c98 100644
+--- a/api/opencl/amdocl/CMakeLists.txt
++++ b/api/opencl/amdocl/CMakeLists.txt
+@@ -21,7 +21,6 @@ include_directories(${CMAKE_SOURCE_DIR}/api/opencl)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos/headers)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos/headers/opencl2.2)
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/driver/src)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib/include)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib/backends/common)
+@@ -67,6 +66,9 @@ add_library(amdocl64 SHARED
+ ${ADDITIONAL_LIBRARIES}
+ )
+
+-target_link_libraries(amdocl64 opencl_driver oclelf pthread dl
${ROCR_LIBRARIES})
++target_link_libraries(amdocl64 opencl_driver oclelf pthread dl
${ROCR_LIBRARIES} LLVMBitWriter LLVMIRReader LLVMLinker LLVMMCParser amd_comgr)
+
+-install(TARGETS amdocl64 LIBRARY DESTINATION lib/x86_64 COMPONENT
applications)
++install(TARGETS amdocl64 LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}
COMPONENT applications)
++
++file(GENERATE OUTPUT "${CMAKE_BINARY_DIR}/amdocl64.icd" CONTENT
"$<TARGET_FILE_NAME:amdocl64>")
++install(FILES "${CMAKE_BINARY_DIR}/amdocl64.icd" DESTINATION
"${CMAKE_INSTALL_FULL_SYSCONFDIR}/OpenCL/vendors/")
+diff --git a/api/opencl/amdocl/cl_execute.cpp
b/api/opencl/amdocl/cl_execute.cpp
+index 0336353..ee799e6 100644
+--- a/api/opencl/amdocl/cl_execute.cpp
++++ b/api/opencl/amdocl/cl_execute.cpp
+@@ -10,7 +10,7 @@
+ #include "platform/program.hpp"
+ #include "os/os.hpp"
+
+-#include <icd/icd_dispatch.h>
++#include <icd/loader/icd_dispatch.h>
+
+ /*! \addtogroup API
+ * @{
+diff --git a/api/opencl/amdocl/cl_icd.cpp b/api/opencl/amdocl/cl_icd.cpp
+index 71b886c..7e4e317 100644
+--- a/api/opencl/amdocl/cl_icd.cpp
++++ b/api/opencl/amdocl/cl_icd.cpp
+@@ -10,7 +10,7 @@
+ #include "cl_d3d11_amd.hpp"
+ #endif //_WIN32
+
+-#include <icd/icd_dispatch.h>
++#include <icd/loader/icd_dispatch.h>
+
+ #include <mutex>
+
+diff --git a/opencl.xml b/opencl.xml
+deleted file mode 100644
+index 4813651..0000000
+--- a/opencl.xml
++++ /dev/null
+@@ -1,17 +0,0 @@
+-<?xml version="1.0" encoding="UTF-8"?>
+-<manifest>
+- <remote name="RadeonOpenCompute"
fetch="https://github.com/RadeonOpenCompute/"/>
+- <remote name="KhronosGroup" fetch="https://github.com/KhronosGroup/"/>
+-
+- <default remote="RadeonOpenCompute" revision="refs/tags/roc-2.6.0"
sync-c="true" sync-j="4"/>
+-
+- <project path="opencl" name="ROCm-OpenCL-Runtime"/>
+- <project path="opencl/compiler/driver" name="ROCm-OpenCL-Driver"/>
+- <project path="opencl/compiler/llvm" name="llvm"
revision="refs/tags/roc-ocl-2.6.0"/>
+- <project path="opencl/compiler/llvm/tools/clang" name="clang"/>
+- <project path="opencl/compiler/llvm/tools/lld" name="lld"
revision="refs/tags/roc-ocl-2.6.0"/>
+-
+- <project path="opencl/library/amdgcn" name="ROCm-Device-Libs"
revision="refs/tags/roc-ocl-2.6.0"/>
+-
+- <project path="opencl/api/opencl/khronos/icd" name="OpenCL-ICD-Loader"
remote="KhronosGroup" revision="7433f2acbf5bbc400f26494ff1dc895da6265bef"/>
+-</manifest>
+diff --git a/runtime/CMakeLists.txt b/runtime/CMakeLists.txt
+index 62e2bf7..5e4abae 100644
+--- a/runtime/CMakeLists.txt
++++ b/runtime/CMakeLists.txt
+@@ -14,8 +14,6 @@ include_directories(${CMAKE_SOURCE_DIR}/api/opencl)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos/headers)
+ include_directories(${CMAKE_SOURCE_DIR}/api/opencl/khronos/headers/opencl2.2)
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/llvm/include)
+-include_directories(${CMAKE_SOURCE_DIR}/compiler/driver/src)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib/include)
+ include_directories(${CMAKE_SOURCE_DIR}/compiler/lib/backends/common)
+@@ -36,6 +34,7 @@ add_library(oclruntime OBJECT
+ utils/flags.cpp
+ utils/debug.cpp
+ device/appprofile.cpp
++ device/comgrctx.cpp
+ device/device.cpp
+ device/hwdebug.cpp
+ device/blitcl.cpp
+@@ -63,6 +62,8 @@ add_library(oclruntime OBJECT
+ ${CMAKE_SOURCE_DIR}/compiler/tools/caching/cache.cpp
+ )
+ set_target_properties(oclruntime PROPERTIES POSITION_INDEPENDENT_CODE ON)
++add_dependencies(oclruntime opencl1.2-c.amdgcn.inc_target)
++add_dependencies(oclruntime opencl2.0-c.amdgcn.inc_target)
+
+ set(AMDGCN_DEP_LIST)
+ if(AMDGCN_TARGETS_LIB_DEPS)
+diff --git a/runtime/device/device.hpp b/runtime/device/device.hpp
+index 450cd69..8134fa2 100644
+--- a/runtime/device/device.hpp
++++ b/runtime/device/device.hpp
+@@ -21,7 +21,7 @@
+
+ #if defined(WITH_LIGHTNING_COMPILER) && !defined(USE_COMGR_LIBRARY)
+ #include "caching/cache.hpp"
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #endif // defined(WITH_LIGHTNING_COMPILER) && ! defined(USE_COMGR_LIBRARY)
+ #include "acl.h"
+
+diff --git a/runtime/device/devprogram.cpp b/runtime/device/devprogram.cpp
+index 22fb125..3d475b2 100644
+--- a/runtime/device/devprogram.cpp
++++ b/runtime/device/devprogram.cpp
+@@ -14,7 +14,7 @@
+
+ #if defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY)
+ #ifndef USE_COMGR_LIBRARY
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #include "libraries.amdgcn.inc"
+ #include "opencl1.2-c.amdgcn.inc"
+ #include "opencl2.0-c.amdgcn.inc"
+diff --git a/runtime/device/devprogram.hpp b/runtime/device/devprogram.hpp
+index 67af239..8390e7d 100644
+--- a/runtime/device/devprogram.hpp
++++ b/runtime/device/devprogram.hpp
+@@ -11,7 +11,7 @@
+
+ #if defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY)
+ #ifndef USE_COMGR_LIBRARY
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #else
+ #include "amd_comgr.h"
+ #endif
+diff --git a/runtime/device/rocm/CMakeLists.txt
b/runtime/device/rocm/CMakeLists.txt
+index 4ce7c45..617bece 100644
+--- a/runtime/device/rocm/CMakeLists.txt
++++ b/runtime/device/rocm/CMakeLists.txt
+@@ -1,3 +1,21 @@
++find_package(Clang REQUIRED CONFIG)
++
++# FIXME: CLANG_CMAKE_DIR seems like the most stable way to find this, but
++# really there is no way to reliably discover this header.
++#
++# We effectively back up to the Clang output directory (for the case of a
build
++# tree) or install prefix (for the case of an installed copy), and then search
++# for a file named opencl-c.h anywhere below that. We take the first result in
++# the case where there are multiple (e.g. if there is an installed copy nested
++# in a build directory). This is a bit imprecise, but it covers cases like
MSVC
++# adding some additional configuration-specific subdirectories to the build
++# tree but not to an installed copy.
++file(GLOB_RECURSE OPENCL_C_H_LIST "${CLANG_CMAKE_DIR}/../../../*/opencl-c.h")
++list(GET OPENCL_C_H_LIST 0 OPENCL_C_H)
++if (NOT EXISTS "${OPENCL_C_H}" OR IS_DIRECTORY "${OPENCL_C_H}")
++ message(FATAL_ERROR "Unable to locate opencl-c.h from the supplied Clang.
The path '${CLANG_CMAKE_DIR}/../../../*' was searched.")
++endif()
++
+ include(bc2h)
+
+ set(INC_SUFFIX "amdgcn.inc")
+@@ -74,8 +92,8 @@ endforeach()
+
+ # generating opencl*.inc files
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/opencl1.2-c.amdgcn.pch
+- COMMAND clang -cc1 -x cl-header -triple amdgcn-amd-amdhsa -Werror -O3
-DNDEBUG -cl-std=CL1.2 -emit-pch -o
${CMAKE_CURRENT_BINARY_DIR}/opencl1.2-c.amdgcn.pch <
${CMAKE_SOURCE_DIR}/compiler/llvm/tools/clang/lib/Headers/opencl-c.h
+- DEPENDS clang
${CMAKE_SOURCE_DIR}/compiler/llvm/tools/clang/lib/Headers/opencl-c.h
++ COMMAND clang -cc1 -x cl-header -triple amdgcn-amd-amdhsa -Werror -O3
-DNDEBUG -cl-std=CL1.2 -emit-pch -o
${CMAKE_CURRENT_BINARY_DIR}/opencl1.2-c.amdgcn.pch < ${OPENCL_C_H}
++ DEPENDS clang ${OPENCL_C_H}
+ COMMENT "Generating opencl1.2-c.amdgcn.pch"
+ )
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/opencl1.2-c.amdgcn.inc
+@@ -88,8 +106,8 @@ add_custom_target(opencl1.2-c.amdgcn.inc_target ALL DEPENDS
${CMAKE_CURRENT_BINA
+ add_dependencies(oclrocm opencl1.2-c.amdgcn.inc_target)
+
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/opencl2.0-c.amdgcn.pch
+- COMMAND clang -cc1 -x cl-header -triple amdgcn-amd-amdhsa -Werror -O3
-DNDEBUG -cl-std=CL2.0 -emit-pch -o
${CMAKE_CURRENT_BINARY_DIR}/opencl2.0-c.amdgcn.pch <
${CMAKE_SOURCE_DIR}/compiler/llvm/tools/clang/lib/Headers/opencl-c.h
+- DEPENDS clang
${CMAKE_SOURCE_DIR}/compiler/llvm/tools/clang/lib/Headers/opencl-c.h
++ COMMAND clang -cc1 -x cl-header -triple amdgcn-amd-amdhsa -Werror -O3
-DNDEBUG -cl-std=CL2.0 -emit-pch -o
${CMAKE_CURRENT_BINARY_DIR}/opencl2.0-c.amdgcn.pch < ${OPENCL_C_H}
++ DEPENDS clang ${OPENCL_C_H}
+ COMMENT "Generating opencl2.0-c.amdgcn.pch"
+ )
+ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/opencl2.0-c.amdgcn.inc
+diff --git a/runtime/device/rocm/rocdevice.cpp
b/runtime/device/rocm/rocdevice.cpp
+index 8b9901b..2648459 100644
+--- a/runtime/device/rocm/rocdevice.cpp
++++ b/runtime/device/rocm/rocdevice.cpp
+@@ -20,7 +20,7 @@
+ #include "device/rocm/rocvirtual.hpp"
+ #include "device/rocm/rocprogram.hpp"
+ #if defined(WITH_LIGHTNING_COMPILER) && ! defined(USE_COMGR_LIBRARY)
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #endif // defined(WITH_LIGHTNING_COMPILER) && ! defined(USE_COMGR_LIBRARY)
+ #include "device/rocm/rocmemory.hpp"
+ #include "device/rocm/rocglinterop.hpp"
+diff --git a/runtime/device/rocm/rockernel.cpp
b/runtime/device/rocm/rockernel.cpp
+index 8a28acc..7d5b62d 100644
+--- a/runtime/device/rocm/rockernel.cpp
++++ b/runtime/device/rocm/rockernel.cpp
+@@ -11,7 +11,7 @@
+
+ #if defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY)
+ #ifndef USE_COMGR_LIBRARY
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #endif
+ #include "llvm/Support/AMDGPUMetadata.h"
+
+diff --git a/runtime/device/rocm/rocprogram.cpp
b/runtime/device/rocm/rocprogram.cpp
+index 191e58c..58d82e6 100644
+--- a/runtime/device/rocm/rocprogram.cpp
++++ b/runtime/device/rocm/rocprogram.cpp
+@@ -11,7 +11,7 @@
+ #include <gelf.h>
+ #include "libraries.amdgcn.inc"
+ #ifndef USE_COMGR_LIBRARY
+-#include "driver/AmdCompiler.h"
++#include "AmdCompiler.h"
+ #endif
+ #endif // defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY)
+
+diff --git a/runtime/platform/object.hpp b/runtime/platform/object.hpp
+index 8ab1b68..b33a9e3 100644
+--- a/runtime/platform/object.hpp
++++ b/runtime/platform/object.hpp
+@@ -9,7 +9,7 @@
+ #include "os/alloc.hpp"
+ #include "thread/monitor.hpp"
+ #include "utils/util.hpp"
+-#include <icd/icd_dispatch.h>
++#include <icd/loader/icd_dispatch.h>
+
+
+ #define KHR_CL_TYPES_DO(F)
\
+diff --git a/tools/clinfo/CMakeLists.txt b/tools/clinfo/CMakeLists.txt
+deleted file mode 100644
+index 974a46f..0000000
+--- a/tools/clinfo/CMakeLists.txt
++++ /dev/null
+@@ -1,13 +0,0 @@
+-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
+-
+-set (CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)
+-
+-include_directories(${OPENCL_INCLUDE_DIRS})
+-
+-add_definitions(-DHAVE_CL2_HPP)
+-
+-add_executable(clinfo clinfo.cpp)
+-
+-target_link_libraries(clinfo OpenCL)
+-
+-install(TARGETS clinfo DESTINATION bin)
+diff --git a/tools/clinfo/clinfo.cpp b/tools/clinfo/clinfo.cpp
+deleted file mode 100644
+index c5afb7f..0000000
+--- a/tools/clinfo/clinfo.cpp
++++ /dev/null
+@@ -1,832 +0,0 @@
+-/* ============================================================
+-
+-Copyright (c) 2010 Advanced Micro Devices, Inc. All rights reserved.
+-
+-Redistribution and use of this material is permitted under the following
+-conditions:
+-
+-Redistributions must retain the above copyright notice and all terms of this
+-license.
+-
+-In no event shall anyone redistributing or accessing or using this material
+-commence or participate in any arbitration or legal action relating to this
+-material against Advanced Micro Devices, Inc. or any copyright holders or
+-contributors. The foregoing shall survive any expiration or termination of
+-this license or any agreement or access or use related to this material.
+-
+-ANY BREACH OF ANY TERM OF THIS LICENSE SHALL RESULT IN THE IMMEDIATE
REVOCATION
+-OF ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE THIS MATERIAL.
+-
+-THIS MATERIAL IS PROVIDED BY ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT
+-HOLDERS AND CONTRIBUTORS "AS IS" IN ITS CURRENT CONDITION AND WITHOUT ANY
+-REPRESENTATIONS, GUARANTEE, OR WARRANTY OF ANY KIND OR IN ANY WAY RELATED TO
+-SUPPORT, INDEMNITY, ERROR FREE OR UNINTERRUPTED OPERA TION, OR THAT IT IS FREE
+-FROM DEFECTS OR VIRUSES. ALL OBLIGATIONS ARE HEREBY DISCLAIMED - WHETHER
+-EXPRESS, IMPLIED, OR STATUTORY - INCLUDING, BUT NOT LIMITED TO, ANY IMPLIED
+-WARRANTIES OF TITLE, MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE,
+-ACCURACY, COMPLETENESS, OPERABILITY, QUALITY OF SERVICE, OR NON-INFRINGEMENT.
+-IN NO EVENT SHALL ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR
+-CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
PUNITIVE,
+-EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
PROCUREMENT
+-OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, REVENUE, DATA, OR PROFITS; OR
+-BUSINESS INTERRUPTION) HOWEVER CAUSED OR BASED ON ANY THEORY OF LIABILITY
+-ARISING IN ANY WAY RELATED TO THIS MATERIAL, EVEN IF ADVISED OF THE
POSSIBILITY
+-OF SUCH DAMAGE. THE ENTIRE AND AGGREGATE LIABILITY OF ADVANCED MICRO DEVICES,
+-INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS SHALL NOT EXCEED TEN DOLLARS
+-(US $10.00). ANYONE REDISTRIBUTING OR ACCESSING OR USING THIS MATERIAL ACCEPTS
+-THIS ALLOCATION OF RISK AND AGREES TO RELEASE ADVANCED MICRO DEVICES, INC. AND
+-ANY COPYRIGHT HOLDERS AND CONTRIBUTORS FROM ANY AND ALL LIABILITIES,
+-OBLIGATIONS, CLAIMS, OR DEMANDS IN EXCESS OF TEN DOLLARS (US $10.00). THE
+-FOREGOING ARE ESSENTIAL TERMS OF THIS LICENSE AND, IF ANY OF THESE TERMS ARE
+-CONSTRUED AS UNENFORCEABLE, FAIL IN ESSENTIAL PURPOSE, OR BECOME VOID OR
+-DETRIMENTAL TO ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR
+-CONTRIBUTORS FOR ANY REASON, THEN ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE
+-THIS MATERIAL SHALL TERMINATE IMMEDIATELY. MOREOVER, THE FOREGOING SHALL
+-SURVIVE ANY EXPIRATION OR TERMINATION OF THIS LICENSE OR ANY AGREEMENT OR
+-ACCESS OR USE RELATED TO THIS MATERIAL.
+-
+-NOTICE IS HEREBY PROVIDED, AND BY REDISTRIBUTING OR ACCESSING OR USING THIS
+-MATERIAL SUCH NOTICE IS ACKNOWLEDGED, THAT THIS MATERIAL MAY BE SUBJECT TO
+-RESTRICTIONS UNDER THE LAWS AND REGULATIONS OF THE UNITED STATES OR OTHER
+-COUNTRIES, WHICH INCLUDE BUT ARE NOT LIMITED TO, U.S. EXPORT CONTROL LAWS SUCH
+-AS THE EXPORT ADMINISTRATION REGULATIONS AND NATIONAL SECURITY CONTROLS AS
+-DEFINED THEREUNDER, AS WELL AS STATE DEPARTMENT CONTROLS UNDER THE U.S.
+-MUNITIONS LIST. THIS MATERIAL MAY NOT BE USED, RELEASED, TRANSFERRED,
IMPORTED,
+-EXPORTED AND/OR RE-EXPORTED IN ANY MANNER PROHIBITED UNDER ANY APPLICABLE
LAWS,
+-INCLUDING U.S. EXPORT CONTROL LAWS REGARDING SPECIFICALLY DESIGNATED PERSONS,
+-COUNTRIES AND NATIONALS OF COUNTRIES SUBJECT TO NATIONAL SECURITY CONTROLS.
+-MOREOVER, THE FOREGOING SHALL SURVIVE ANY EXPIRATION OR TERMINATION OF ANY
+-LICENSE OR AGREEMENT OR ACCESS OR USE RELATED TO THIS MATERIAL.
+-
+-NOTICE REGARDING THE U.S. GOVERNMENT AND DOD AGENCIES: This material is
+-provided with "RESTRICTED RIGHTS" and/or "LIMITED RIGHTS" as applicable to
+-computer software and technical data, respectively. Use, duplication,
+-distribution or disclosure by the U.S. Government and/or DOD agencies is
+-subject to the full extent of restrictions in all applicable regulations,
+-including those found at FAR52.227 and DFARS252.227 et seq. and any successor
+-regulations thereof. Use of this material by the U.S. Government and/or DOD
+-agencies is acknowledgment of the proprietary rights of any copyright holders
+-and contributors, including those of Advanced Micro Devices, Inc., as well as
+-the provisions of FAR52.227-14 through 23 regarding privately developed and/or
+-commercial computer software.
+-
+-This license forms the entire agreement regarding the subject matter hereof
and
+-supersedes all proposals and prior discussions and writings between the
parties
+-with respect thereto. This license does not affect any ownership, rights,
title,
+-or interest in, or relating to, this material. No terms of this license can be
+-modified or waived, and no breach of this license can be excused, unless done
+-so in a writing signed by all affected parties. Each term of this license is
+-separately enforceable. If any term of this license is determined to be or
+-becomes unenforceable or illegal, such term shall be reformed to the minimum
+-extent necessary in order for this license to remain in effect in accordance
+-with its terms as modified by such reformation. This license shall be governed
+-by and construed in accordance with the laws of the State of Texas without
+-regard to rules on conflicts of law of any state or jurisdiction or the United
+-Nations Convention on the International Sale of Goods. All disputes arising
out
+-of this license shall be subject to the jurisdiction of the federal and state
+-courts in Austin, Texas, and all defenses are hereby waived concerning
personal
+-jurisdiction and venue of these courts.
+-
+-============================================================ */
+-
+-#include <iostream>
+-#include <map>
+-#include <set>
+-#include <algorithm>
+-#include <string>
+-#include <cstdlib>
+-#include <cstdio>
+-#if !defined(_WIN32)
+-#include <errno.h>
+-#endif
+-
+-#ifdef _MSC_VER
+-#pragma warning(disable: 4290)
+-#endif
+-
+-#if defined(HAVE_CL2_HPP)
+-#define CL_HPP_ENABLE_EXCEPTIONS
+-#define CL_HPP_MINIMUM_OPENCL_VERSION 120
+-#define CL_HPP_TARGET_OPENCL_VERSION 200
+-#define CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY
+-#include "CL/cl2.hpp"
+-#else // !HAVE_CL2_HPP
+-#define __CL_ENABLE_EXCEPTIONS
+-#define __MAX_DEFAULT_VECTOR_SIZE 50
+-#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
+-#define CL_USE_DEPRECATED_OPENCL_2_0_APIS
+-#include "cl.hpp"
+-#endif // !HAVE_CL2_HPP
+-
+-bool verbose = false;
+-
+-/// Returns EXIT_SUCCESS on success, EXIT_FAILURE on failure.
+-int
+-main(int argc, char** argv)
+-{
+- /* Error flag */
+- cl_int err;
+-
+- //parse input
+- for(int i = 1; i < argc; i++){
+- if ((strcmp(argv[i], "-v") == 0) ||
+- (strcmp(argv[i], "--verbose") == 0)){
+- verbose = true;
+- } else if ((strcmp(argv[i], "-h") == 0) ||
+- (strcmp(argv[i], "--help") == 0)){
+- std::cout << "Usage is: " << argv[0] << " [-v|--verbose]" <<
std::endl;
+- return EXIT_FAILURE;
+- }
+- }
+-
+- // Platform info
+- std::vector<cl::Platform> platforms;
+-
+- try {
+- err = cl::Platform::get(&platforms);
+-
+- // Iteratate over platforms
+- std::cout << "Number of platforms:\t\t\t\t "
+- << platforms.size()
+- << std::endl;
+- for (std::vector<cl::Platform>::iterator i = platforms.begin();
+- i != platforms.end();
+- ++i) {
+- const cl::Platform& platform = *i;
+-
+- std::cout << " Platform Profile:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_PROFILE>().c_str()
+- << std::endl;
+- std::cout << " Platform Version:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_VERSION>().c_str()
+- << std::endl;
+- std::cout << " Platform Name:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_NAME>().c_str()
+- << std::endl;
+- std::cout << " Platform Vendor:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_VENDOR>().c_str() <<
std::endl;
+- if (platform.getInfo<CL_PLATFORM_EXTENSIONS>().size() > 0) {
+- std::cout << " Platform Extensions:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_EXTENSIONS>().c_str()
+- << std::endl;
+- }
+- }
+-
+- std::cout << std::endl << std:: endl;
+- // Now Iteratate over each platform and its devices
+- for (std::vector<cl::Platform>::iterator p = platforms.begin();
+- p != platforms.end();
+- ++p) {
+- const cl::Platform& platform = *p;
+- std::cout << " Platform Name:\t\t\t\t "
+- << platform.getInfo<CL_PLATFORM_NAME>().c_str()
+- << std::endl;
+-
+- std::vector<cl::Device> devices;
+- platform.getDevices(CL_DEVICE_TYPE_ALL, &devices);
+-
+- // Get OpenCL version
+- std::string platformVersionStr =
platform.getInfo<CL_PLATFORM_VERSION>();
+- std::string openclVerstionStr(platformVersionStr.c_str());
+- size_t vStart = openclVerstionStr.find(" ", 0);
+- size_t vEnd = openclVerstionStr.find(" ", vStart + 1);
+- std::string vStrVal = openclVerstionStr.substr(vStart + 1, vEnd -
vStart - 1);
+-
+- std::cout << "Number of devices:\t\t\t\t " << devices.size() <<
std::endl;
+- for (std::vector<cl::Device>::iterator i = devices.begin();
+- i != devices.end();
+- ++i) {
+- const cl::Device& device = *i;
+- /* Get device name */
+- std::string deviceName = device.getInfo<CL_DEVICE_NAME>();
+- cl_device_type dtype = device.getInfo<CL_DEVICE_TYPE>();
+-
+- /* Get CAL driver version in int */
+- std::string driverVersion = device.getInfo<CL_DRIVER_VERSION>();
+- std::string calVersion(driverVersion.c_str());
+- calVersion = calVersion.substr(calVersion.find_last_of(".") + 1);
+- int version = atoi(calVersion.c_str());
+-
+- std::cout << " Device Type:\t\t\t\t\t " ;
+- switch (dtype) {
+- case CL_DEVICE_TYPE_ACCELERATOR:
+- std::cout << "CL_DEVICE_TYPE_ACCRLERATOR" << std::endl;
+- break;
+- case CL_DEVICE_TYPE_CPU:
+- std::cout << "CL_DEVICE_TYPE_CPU" << std::endl;
+- break;
+- case CL_DEVICE_TYPE_DEFAULT:
+- std::cout << "CL_DEVICE_TYPE_DEFAULT" << std::endl;
+- break;
+- case CL_DEVICE_TYPE_GPU:
+- std::cout << "CL_DEVICE_TYPE_GPU" << std::endl;
+- break;
+- }
+-
+- std::cout << " Vendor ID:\t\t\t\t\t "
+- << std::hex
+- << device.getInfo<CL_DEVICE_VENDOR_ID>()
+- << "h"
+- << std::dec
+- << std::endl;
+-
+- bool isAMDPlatform =
(strcmp(platform.getInfo<CL_PLATFORM_NAME>().c_str(), "AMD Accelerated Parallel
Processing") == 0) ? true : false;
+- if (isAMDPlatform)
+- {
+- std::string boardName;
+- device.getInfo(CL_DEVICE_BOARD_NAME_AMD, &boardName);
+- std::cout << " Board name:\t\t\t\t\t "
+- << boardName.c_str()
+- << std::endl;
+-
+- cl_device_topology_amd topology;
+- err = device.getInfo(CL_DEVICE_TOPOLOGY_AMD, &topology);
+- if (topology.raw.type == CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD) {
+- std::cout << " Device Topology:\t\t\t\t "
+- << "PCI[ B#" << (int)topology.pcie.bus
+- << ", D#" << (int)topology.pcie.device
+- << ", F#" << (int)topology.pcie.function
+- << " ]" << std::endl;
+- }
+- }
+-
+- std::cout << " Max compute units:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>()
+- << std::endl;
+-
+- std::cout << " Max work items dimensions:\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS>()
+- << std::endl;
+-
+- std::vector< ::size_t> witems =
+- device.getInfo<CL_DEVICE_MAX_WORK_ITEM_SIZES>();
+- for (unsigned int x = 0;
+- x < device.getInfo<CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS>();
+- x++) {
+- std::cout << " Max work items["
+- << x << "]:\t\t\t\t "
+- << witems[x]
+- << std::endl;
+- }
+-
+- std::cout << " Max work group size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_WORK_GROUP_SIZE>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width char:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width short:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width int:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width long:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width float:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT>()
+- << std::endl;
+-
+- std::cout << " Preferred vector width double:\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE>()
+- << std::endl;
+-
+-#ifdef CL_VERSION_1_1
+- if(vStrVal.compare("1.0") > 0)
+- {
+- std::cout << " Native vector width char:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR>()
+- << std::endl;
+-
+- std::cout << " Native vector width short:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT>()
+- << std::endl;
+-
+- std::cout << " Native vector width int:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_INT>()
+- << std::endl;
+-
+- std::cout << " Native vector width long:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG>()
+- << std::endl;
+-
+- std::cout << " Native vector width float:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT>()
+- << std::endl;
+-
+- std::cout << " Native vector width double:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE>()
+- << std::endl;
+- }
+-#endif // CL_VERSION_1_1
+- std::cout << " Max clock frequency:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_CLOCK_FREQUENCY>()
+- << "Mhz"
+- << std::endl;
+-
+- std::cout << " Address bits:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_ADDRESS_BITS>()
+- << std::endl;
+-
+- std::cout << " Max memory allocation:\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_MEM_ALLOC_SIZE>()
+- << std::endl;
+-
+- std::cout << " Image support:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_IMAGE_SUPPORT>() ? "Yes" :
"No")
+- << std::endl;
+-
+- if (device.getInfo<CL_DEVICE_IMAGE_SUPPORT>())
+- {
+-
+- std::cout << " Max number of images read arguments:\t\t "
+- << device.getInfo<CL_DEVICE_MAX_READ_IMAGE_ARGS>()
+- << std::endl;
+-
+- std::cout << " Max number of images write arguments:\t\t "
+- << device.getInfo<CL_DEVICE_MAX_WRITE_IMAGE_ARGS>()
+- << std::endl;
+-
+- std::cout << " Max image 2D width:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_IMAGE2D_MAX_WIDTH>()
+- << std::endl;
+-
+- std::cout << " Max image 2D height:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_IMAGE2D_MAX_HEIGHT>()
+- << std::endl;
+-
+- std::cout << " Max image 3D width:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_IMAGE3D_MAX_WIDTH>()
+- << std::endl;
+-
+- std::cout << " Max image 3D height:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_IMAGE3D_MAX_HEIGHT>()
+- << std::endl;
+-
+- std::cout << " Max image 3D depth:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_IMAGE3D_MAX_DEPTH>()
+- << std::endl;
+-
+- std::cout << " Max samplers within kernel:\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_SAMPLERS>()
+- << std::endl;
+-
+- if (verbose)
+- {
+- std::cout << " Image formats supported:" << std::endl;
+- std::vector<cl::ImageFormat> formats;
+-
+- cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM,
(cl_context_properties)(*p)(), 0 };
+- std::vector<cl::Device> device;
+- device.push_back(*i);
+- cl::Context context(device, cps, NULL, NULL, &err);
+-
+- std::map<int,std::string> channelOrder;
+- channelOrder[CL_R] = "CL_R";
+- channelOrder[CL_A] = "CL_A";
+- channelOrder[CL_RG] = "CL_RG";
+- channelOrder[CL_RA] = "CL_RA";
+- channelOrder[CL_RGB] = "CL_RGB";
+- channelOrder[CL_RGBA] = "CL_RGBA";
+- channelOrder[CL_BGRA] = "CL_BGRA";
+- channelOrder[CL_ARGB] = "CL_ARGB";
+- channelOrder[CL_INTENSITY] = "CL_INTENSITY";
+- channelOrder[CL_LUMINANCE] = "CL_LUMINANCE";
+- channelOrder[CL_Rx] = "CL_Rx";
+- channelOrder[CL_RGx] = "CL_RGx";
+- channelOrder[CL_RGBx] = "CL_RGBx";
+-
+- std::map<int,std::pair<std::string, std::string> >
channelType;
+- channelType[CL_SNORM_INT8] = std::make_pair("snorm",
"int8");
+- channelType[CL_SNORM_INT16] = std::make_pair("snorm",
"int16");
+- channelType[CL_UNORM_INT8] = std::make_pair("unorm",
"int8");
+- channelType[CL_UNORM_INT16] = std::make_pair("unorm",
"int16");
+- channelType[CL_UNORM_SHORT_565] = std::make_pair("unorm",
"short_565");
+- channelType[CL_UNORM_SHORT_555] = std::make_pair("unorm",
"short_555");
+- channelType[CL_UNORM_INT_101010] =
std::make_pair("unorm", "int_101010");
+- channelType[CL_SIGNED_INT8] = std::make_pair("signed",
"int8");
+- channelType[CL_SIGNED_INT16] = std::make_pair("signed",
"int16");
+- channelType[CL_SIGNED_INT32] = std::make_pair("signed",
"int32");
+- channelType[CL_UNSIGNED_INT8] =
std::make_pair("unsigned", "int8");
+- channelType[CL_UNSIGNED_INT16] =
std::make_pair("unsigned", "int16");
+- channelType[CL_UNSIGNED_INT32] =
std::make_pair("unsigned", "int32");
+- channelType[CL_HALF_FLOAT] = std::make_pair("half_float",
"");
+- channelType[CL_FLOAT] = std::make_pair("float", "");
+-
+- std::vector<std::pair<int, std::string> > imageDimensions;
+-
imageDimensions.push_back(std::make_pair(CL_MEM_OBJECT_IMAGE2D, std::string("2D
")));
+-
imageDimensions.push_back(std::make_pair(CL_MEM_OBJECT_IMAGE3D, std::string("3D
")));
+- for(std::vector<std::pair<int, std::string> >::iterator
id = imageDimensions.begin();
+- id != imageDimensions.end();
+- id++){
+-
+- struct imageAccessStruct {
+- std::string name;
+- int access;
+- std::vector<cl::ImageFormat> formats;
+- } imageAccess[] =
{{std::string("Read-Write/Read-Only/Write-Only"), CL_MEM_READ_WRITE,
std::vector<cl::ImageFormat>()},
+- {std::string("Read-Only"),
CL_MEM_READ_ONLY, std::vector<cl::ImageFormat>()},
+- {std::string("Write-Only"),
CL_MEM_WRITE_ONLY, std::vector<cl::ImageFormat>()}};
+-
+- for(size_t ia=0; ia <
sizeof(imageAccess)/sizeof(imageAccessStruct); ia++){
+-
context.getSupportedImageFormats(imageAccess[ia].access, (*id).first,
&(imageAccess[ia].formats));
+- bool printTopHeader = true;
+- for (std::map<int,std::string>::iterator o =
channelOrder.begin();
+- o != channelOrder.end();
+- o++)
+- {
+- bool printHeader = true;
+-
+- for (std::vector<cl::ImageFormat>::iterator
it = imageAccess[ia].formats.begin();
+- it != imageAccess[ia].formats.end();
+- ++it)
+- {
+- if ( (*o).first ==
(int)(*it).image_channel_order)
+- {
+- bool printedAlready = false;
+- //see if this was already print in
RW/RO/WO
+- if (ia !=0)
+- {
+- for
(std::vector<cl::ImageFormat>::iterator searchIt =
imageAccess[0].formats.begin();
+- searchIt !=
imageAccess[0].formats.end();
+- searchIt++)
+- {
+- if (
((*searchIt).image_channel_data_type == (*it).image_channel_data_type) &&
+-
((*searchIt).image_channel_order == (*it).image_channel_order))
+- {
+- printedAlready = true;
+- break;
+- }
+- }
+- }
+- if (printedAlready)
+- {
+- continue;
+- }
+- if (printTopHeader)
+- {
+- std::cout << " " <<
(*id).second << imageAccess[ia].name << std::endl;
+- printTopHeader = false;
+- }
+- if (printHeader)
+- {
+- std::cout << " " <<
(*o).second << ": ";
+- printHeader = false;
+- }
+- std::cout <<
channelType[(*it).image_channel_data_type].first;
+- if
(channelType[(*it).image_channel_data_type].second != "")
+- {
+- std::cout << "-"
+- <<
channelType[(*it).image_channel_data_type].second;
+- }
+- if (it !=
(imageAccess[ia].formats.end() - 1))
+- {
+- std::cout << " ";
+- }
+- }
+- }
+- if (printHeader == false)
+- {
+- std::cout << std::endl;
+- }
+- }
+- }
+- }
+- }
+- }
+-
+- std::cout << " Max size of kernel argument:\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_PARAMETER_SIZE>()
+- << std::endl;
+-
+- std::cout << " Alignment (bits) of base address:\t\t "
+- << device.getInfo<CL_DEVICE_MEM_BASE_ADDR_ALIGN>()
+- << std::endl;
+-
+- std::cout << " Minimum alignment (bytes) for any datatype:\t "
+- << device.getInfo<CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE>()
+- << std::endl;
+-
+- std::cout << " Single precision floating point capability" <<
std::endl;
+- std::cout << " Denorms:\t\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_DENORM ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Quiet NaNs:\t\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_INF_NAN ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Round to nearest even:\t\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_ROUND_TO_NEAREST ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Round to zero:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_ROUND_TO_ZERO ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Round to +ve and infinity:\t\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_ROUND_TO_INF ? "Yes" : "No")
+- << std::endl;
+- std::cout << " IEEE754-2008 fused multiply-add:\t\t "
+- << (device.getInfo<CL_DEVICE_SINGLE_FP_CONFIG>() &
+- CL_FP_FMA ? "Yes" : "No")
+- << std::endl;
+-
+- std::cout << " Cache type:\t\t\t\t\t " ;
+- switch (device.getInfo<CL_DEVICE_GLOBAL_MEM_CACHE_TYPE>()) {
+- case CL_NONE:
+- std::cout << "None" << std::endl;
+- break;
+- case CL_READ_ONLY_CACHE:
+- std::cout << "Read only" << std::endl;
+- break;
+- case CL_READ_WRITE_CACHE:
+- std::cout << "Read/Write" << std::endl;
+- break;
+- }
+-
+- std::cout << " Cache line size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE>()
+- << std::endl;
+-
+- std::cout << " Cache size:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_GLOBAL_MEM_CACHE_SIZE>()
+- << std::endl;
+-
+- std::cout << " Global memory size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>()
+- << std::endl;
+-
+- std::cout << " Constant buffer size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE>()
+- << std::endl;
+-
+- std::cout << " Max number of constant args:\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_CONSTANT_ARGS>()
+- << std::endl;
+-
+- std::cout << " Local memory type:\t\t\t\t " ;
+- switch (device.getInfo<CL_DEVICE_LOCAL_MEM_TYPE>()) {
+- case CL_LOCAL:
+- std::cout << "Scratchpad" << std::endl;
+- break;
+- case CL_GLOBAL:
+- std::cout << "Global" << std::endl;
+- break;
+- }
+-
+-
+- std::cout << " Local memory size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>()
+- << std::endl;
+-
+-#if defined(CL_VERSION_2_0)
+- if(vStrVal.compare("2") > 0)
+- {
+- std::cout << " Max pipe arguments:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_PIPE_ARGS>()
+- << std::endl;
+-
+- std::cout << " Max pipe active reservations:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS>()
+- << std::endl;
+-
+- std::cout << " Max pipe packet size:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_PIPE_MAX_PACKET_SIZE>()
+- << std::endl;
+-
+- std::cout << " Max global variable size:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE>()
+- << std::endl;
+-
+- std::cout << " Max global variable preferred total size:\t "
+- <<
device.getInfo<CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE>()
+- << std::endl;
+-
+- std::cout << " Max read/write image args:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS>()
+- << std::endl;
+-
+- std::cout << " Max on device events:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_ON_DEVICE_EVENTS>()
+- << std::endl;
+-
+- std::cout << " Queue on device max size:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE>()
+- << std::endl;
+-
+- std::cout << " Max on device queues:\t\t\t\t "
+- << device.getInfo<CL_DEVICE_MAX_ON_DEVICE_QUEUES>()
+- << std::endl;
+-
+- std::cout << " Queue on device preferred size:\t\t "
+- <<
device.getInfo<CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE>()
+- << std::endl;
+-
+- std::cout << " SVM capabilities:\t\t\t\t " << std::endl;
+- std::cout << " Coarse grain buffer:\t\t\t "
+- << (device.getInfo<CL_DEVICE_SVM_CAPABILITIES>() &
+- CL_DEVICE_SVM_COARSE_GRAIN_BUFFER ? "Yes" :
"No")
+- << std::endl;
+- std::cout << " Fine grain buffer:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SVM_CAPABILITIES>() &
+- CL_DEVICE_SVM_FINE_GRAIN_BUFFER ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Fine grain system:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SVM_CAPABILITIES>() &
+- CL_DEVICE_SVM_FINE_GRAIN_SYSTEM ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Atomics:\t\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_SVM_CAPABILITIES>() &
+- CL_DEVICE_SVM_ATOMICS ? "Yes" : "No")
+- << std::endl;
+-
+- std::cout << " Preferred platform atomic alignment:\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT>()
+- << std::endl;
+-
+- std::cout << " Preferred global atomic alignment:\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT>()
+- << std::endl;
+-
+- std::cout << " Preferred local atomic alignment:\t\t "
+- <<
device.getInfo<CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT>()
+- << std::endl;
+- }
+-#endif // CL_VERSION_2_0
+-
+-#if defined(CL_VERSION_1_1) && !defined(ATI_ARCH_ARM)
+- if(vStrVal.compare("1.0") > 0)
+- {
+- cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM,
(cl_context_properties)(*p)(), 0 };
+-
+- std::vector<cl::Device> device;
+- device.push_back(*i);
+-
+- cl::Context context(device, cps, NULL, NULL, &err);
+- if (err != CL_SUCCESS) {
+- std::cerr << "Context::Context() failed (" << err <<
")\n";
+- return EXIT_FAILURE;
+- }
+- std::string kernelStr("__kernel void hello(){ size_t i =
get_global_id(0); size_t j = get_global_id(1);}");
+- cl::Program::Sources sources(1,
std::make_pair(kernelStr.data(), kernelStr.size()));
+-
+- cl::Program program = cl::Program(context, sources, &err);
+- if (err != CL_SUCCESS) {
+- std::cerr << "Program::Program() failed (" << err <<
")\n";
+- return EXIT_FAILURE;
+- }
+-
+- err = program.build(device);
+- if (err != CL_SUCCESS) {
+-
+- if(err == CL_BUILD_PROGRAM_FAILURE)
+- {
+- std::string str =
program.getBuildInfo<CL_PROGRAM_BUILD_LOG>((*i));
+-
+- std::cout << " \n\t\t\tBUILD LOG\n";
+- std::cout << "
************************************************\n";
+- std::cout << str.c_str() << std::endl;
+- std::cout << "
************************************************\n";
+- }
+-
+- std::cerr << "Program::build() failed (" << err << ")\n";
+- return EXIT_FAILURE;
+- }
+-
+- cl::Kernel kernel(program, "hello", &err);
+- if (err != CL_SUCCESS) {
+- std::cerr << "Kernel::Kernel() failed (" << err << ")\n";
+- return EXIT_FAILURE;
+- }
+-
+- std::cout << " Kernel Preferred work group size multiple:\t "
+- <<
kernel.getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>((*i),
&err)
+- << std::endl;
+- }
+-
+-#endif // CL_VERSION_1_1
+-
+- std::cout << " Error correction support:\t\t\t "
+- << device.getInfo<CL_DEVICE_ERROR_CORRECTION_SUPPORT>()
+- << std::endl;
+-#ifdef CL_VERSION_1_1
+- if(vStrVal.compare("1.0") > 0)
+- {
+- std::cout << " Unified memory for Host and Device:\t\t "
+- << device.getInfo<CL_DEVICE_HOST_UNIFIED_MEMORY>()
+- << std::endl;
+- }
+-#endif // CL_VERSION_1_1
+- std::cout << " Profiling timer resolution:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_PROFILING_TIMER_RESOLUTION>()
+- << std::endl;
+-
+- std::cout << " Device endianess:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_ENDIAN_LITTLE>() ?
"Little" : "Big")
+- << std::endl;
+-
+- std::cout << " Available:\t\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_AVAILABLE>() ? "Yes" :
"No")
+- << std::endl;
+-
+- std::cout << " Compiler available:\t\t\t\t "
+- << (device.getInfo<CL_DEVICE_COMPILER_AVAILABLE>() ?
"Yes" : "No")
+- << std::endl;
+-
+- std::cout << " Execution capabilities:\t\t\t\t " << std::endl;
+- std::cout << " Execute OpenCL kernels:\t\t\t "
+- << (device.getInfo<CL_DEVICE_EXECUTION_CAPABILITIES>() &
+- CL_EXEC_KERNEL ? "Yes" : "No")
+- << std::endl;
+- std::cout << " Execute native function:\t\t\t "
+- << (device.getInfo<CL_DEVICE_EXECUTION_CAPABILITIES>() &
+- CL_EXEC_NATIVE_KERNEL ? "Yes" : "No")
+- << std::endl;
+-
+- std::cout << " Queue on Host properties:\t\t\t\t " << std::endl;
+- std::cout << " Out-of-Order:\t\t\t\t "
+- <<
(device.getInfo<CL_DEVICE_QUEUE_ON_HOST_PROPERTIES>() &
+- CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ? "Yes" :
"No")
+- << std::endl;
+- std::cout << " Profiling :\t\t\t\t\t "
+- <<
(device.getInfo<CL_DEVICE_QUEUE_ON_HOST_PROPERTIES>() &
+- CL_QUEUE_PROFILING_ENABLE ? "Yes" : "No")
+- << std::endl;
+-
+-#ifdef CL_VERSION_2_0
+- if(vStrVal.compare("2") > 0)
+- {
+- std::cout << " Queue on Device properties:\t\t\t\t " <<
std::endl;
+- std::cout << " Out-of-Order:\t\t\t\t "
+- <<
(device.getInfo<CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES>() &
+- CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ? "Yes"
: "No")
+- << std::endl;
+- std::cout << " Profiling :\t\t\t\t\t "
+- <<
(device.getInfo<CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES>() &
+- CL_QUEUE_PROFILING_ENABLE ? "Yes" : "No")
+- << std::endl;
+- }
+-#endif
+-
+- std::cout << " Platform ID:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_PLATFORM>()
+- << std::endl;
+-
+- std::cout << " Name:\t\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_NAME>().c_str()
+- << std::endl;
+-
+- std::cout << " Vendor:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_VENDOR>().c_str()
+- << std::endl;
+-#ifdef CL_VERSION_1_1
+- if(vStrVal.compare("1.0") > 0)
+- {
+- std::cout << " Device OpenCL C version:\t\t\t "
+- <<
device.getInfo<CL_DEVICE_OPENCL_C_VERSION>().c_str()
+- << std::endl;
+- }
+-#endif // CL_VERSION_1_1
+- std::cout << " Driver version:\t\t\t\t "
+- << device.getInfo<CL_DRIVER_VERSION>().c_str()
+- << std::endl;
+-
+- std::cout << " Profile:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_PROFILE>().c_str()
+- << std::endl;
+-
+- std::cout << " Version:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_VERSION>().c_str()
+- << std::endl;
+-
+-
+- std::cout << " Extensions:\t\t\t\t\t "
+- << device.getInfo<CL_DEVICE_EXTENSIONS>().c_str()
+- << std::endl;
+-
+- std::cout << std::endl << std::endl;
+- }
+- }
+- }
+- catch (cl::Error err)
+- {
+- std::cerr
+- << "ERROR: "
+- << err.what()
+- << "("
+- << err.err()
+- << ")"
+- << std::endl;
+- return EXIT_FAILURE;
+- }
+-
+- return EXIT_SUCCESS;
+-}
diff --git a/dev-libs/rocm-opencl-runtime/metadata.xml
b/dev-libs/rocm-opencl-runtime/metadata.xml
new file mode 100644
index 00000000000..359755b2e5a
--- /dev/null
+++ b/dev-libs/rocm-opencl-runtime/metadata.xml
@@ -0,0 +1,14 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<!DOCTYPE pkgmetadata SYSTEM "http://www.gentoo.org/dtd/metadata.dtd";>
+<pkgmetadata>
+ <maintainer type="person">
+ <email>[email protected]</email>
+ <name>Craig Andrews</name>
+ </maintainer>
+ <upstream>
+ <remote-id
type="github">RadeonOpenCompute/ROCm-OpenCL-Runtime</remote-id>
+ </upstream>
+ <longdescription lang="en">
+ Radeon Open Compute OpenCL Compatible Runtime
+ </longdescription>
+</pkgmetadata>
diff --git a/dev-libs/rocm-opencl-runtime/rocm-opencl-runtime-2.6.0.ebuild
b/dev-libs/rocm-opencl-runtime/rocm-opencl-runtime-2.6.0.ebuild
new file mode 100644
index 00000000000..cd44ed80e4d
--- /dev/null
+++ b/dev-libs/rocm-opencl-runtime/rocm-opencl-runtime-2.6.0.ebuild
@@ -0,0 +1,50 @@
+# Copyright 1999-2019 Gentoo Authors
+# Distributed under the terms of the GNU General Public License v2
+
+EAPI=7
+
+inherit cmake-utils
+
+OPENCL_ICD_COMMIT="bc9728edf8cace79cf33bf75560be88fc2432dc4"
+SRC_URI="https://github.com/KhronosGroup/OpenCL-ICD-Loader/archive/${OPENCL_ICD_COMMIT}.tar.gz
-> OpenCL-ICD-Loader-${OPENCL_ICD_COMMIT}.tar.gz"
+if [[ ${PV} == *9999 ]] ; then
+
EGIT_REPO_URI="https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime/";
+ inherit git-r3
+else
+ SRC_URI+="
https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime/archive/roc-${PV}.tar.gz
-> ${P}.tar.gz"
+ KEYWORDS="~amd64"
+ S="${WORKDIR}/ROCm-OpenCL-Runtime-roc-${PV}"
+fi
+
+DESCRIPTION="Radeon Open Compute OpenCL Compatible Runtime"
+HOMEPAGE="https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime";
+
+LICENSE="Apache-2.0 MIT"
+SLOT="0/$(ver_cut 1-2)"
+
+RDEPEND="dev-libs/rocr-runtime
+ dev-libs/rocm-comgr
+ dev-libs/rocm-device-libs
+ dev-libs/rocm-opencl-driver"
+DEPEND="${RDEPEND}
+ dev-lang/ocaml
+ dev-ml/findlib"
+
+PATCHES=(
+ "${FILESDIR}/${P}-unbundle-dependencies.patch"
+)
+
+src_prepare() {
+ mkdir -p "${S}"/api/opencl/khronos/ || die
+ mv "${WORKDIR}/OpenCL-ICD-Loader-${OPENCL_ICD_COMMIT}"
"${S}"/api/opencl/khronos/icd || die
+ cmake-utils_src_prepare
+}
+
+src_configure() {
+ local mycmakeargs=(
+ -DLLVM_DIR="${EPREFIX}/usr/lib/llvm/roc/"
+ -DClang_DIR="${EPREFIX}/usr/lib/llvm/roc/lib/cmake/clang/"
+ -DCMAKE_INSTALL_PREFIX="${EPREFIX}/usr/"
+ )
+ cmake-utils_src_configure
+}
