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;
-}
