From 2cab8307b96534865e73a4470bed122336b86748 Mon Sep 17 00:00:00 2001 From: Ramesh Errabolu Date: Mon, 2 Oct 2017 13:35:14 -0500 Subject: [PATCH] Cmake project to build Rocm Code Objects Change-Id: If3a631615316c203318bb5ae1df328a66e2919b1 --- rocrtst/Kernels/CMakeLists.txt | 162 ++++++++++++++++++++++++ rocrtst/Kernels/binary_search_kernel.cl | 127 +++++++++++++++++++ 2 files changed, 289 insertions(+) create mode 100644 rocrtst/Kernels/CMakeLists.txt create mode 100755 rocrtst/Kernels/binary_search_kernel.cl diff --git a/rocrtst/Kernels/CMakeLists.txt b/rocrtst/Kernels/CMakeLists.txt new file mode 100644 index 0000000000..c2f7506ca6 --- /dev/null +++ b/rocrtst/Kernels/CMakeLists.txt @@ -0,0 +1,162 @@ + +cmake_minimum_required(VERSION 2.8.0) + +# +# Setup build environment +# +# 1) Setup env var LLVM_DIR and OCL_BITCODE_DIR to point to +# folders containing relevant libraries seperately +# +# export LLVM_DIR="Path to Lightning build artifacts" +# +# export OCL_BITCODE_DIR="Path containing Opencl Bitcode libraries" +# +# 2) Make an new folder called build under root folder +# +# mkdir build +# +# 3) Enter into folder of build, and run CMAKE to generate makefile +# and make it +# +# cd build; cmake ..; make +# + +if(WIN32) + message("Windows platform is not supported") + return() +endif() + +# +# Flag to enable / disable verbose output. +# +SET( CMAKE_VERBOSE_MAKEFILE on ) + +set(PROJECT_NAME "CompileKernels") +project (${PROJECT_NAME}) + +# +# Validate LLVM related resources are available +# +if (NOT DEFINED ENV{LLVM_DIR}) + message("LLVM_DIR define is not set. Kernels cannot be built.") + return() +endif() + +# +# Validate Opencl related resources are available +# +if (NOT DEFINED ENV{OCL_BITCODE_DIR}) + message(FATAL_ERROR "OCL_BITCODE_DIR define is not set. Kernels cannot be built.") +endif() + +set(CLANG $ENV{LLVM_DIR}/clang) +if (NOT EXISTS ${CLANG}) + message("Path to clang (${CLANG}) is not valid. Is LLVM_DIR defined correctly?") + return() +endif() + +# +# Define Opencl version if it is not defined +# +if (DEFINED ENV{OPENCL_VER}) + set(OPENCL_VER $ENV{OPENCL_VER}) +else() + message("OPENCL_VER define is not set. Using default") + set(OPENCL_VER "2.0") +endif() + +# +# Define list of Target Device types for which to get code objects +# +set(DEV_LIST "gfx803" "gfx900" CACHE STRING "List of Gfx Devices") +set(TARGET_DEV_LIST ${DEV_LIST}) +separate_arguments(TARGET_DEV_LIST) + +# Maintains a global list of targets to build +set (ROCM_CODEOBJ_LIST "" CACHE INTERNAL ROCM_CODEOBJ_LIST) + +# +# Options that are passed along to Clang to enable code object generation +# +set(KERN_SUFFIX "kernels.hsaco") +set(BITCODE_PREF "-Xclang -mlink-bitcode-file -Xclang") +set(COMMON_BITCODE_LIB1 "${BITCODE_PREF} $ENV{OCL_BITCODE_DIR}/opencl.amdgcn.bc") +set(COMMON_BITCODE_LIB2 "${BITCODE_PREF} $ENV{OCL_BITCODE_DIR}/ockl.amdgcn.bc") +set(COMMON_BITCODE_LIB3 "${BITCODE_PREF} $ENV{OCL_BITCODE_DIR}/ocml.amdgcn.bc") + +# +# Compiles Opencl kernel into a AMDGcn code object +# +function(CompileKernel KRNL_NAME TARGET_DEV) + + # + # Bind names for code object file and directory containing it + set(KERNEL_DIR ${PROJECT_BINARY_DIR}/${TARGET_DEV}) + set(CODEASM_FILE "${KRNL_NAME}_${TARGET_DEV}.asm") + set(CODEOBJ_FILE "${KRNL_NAME}_${TARGET_DEV}.hsaco") + + # + # Add target name to a global list of target names. This must + # be executed before the add_custom_target rule + set(TARGET_NAME "${KRNL_NAME}_${TARGET_DEV}") + set(ROCM_CODEOBJ_LIST ${ROCM_CODEOBJ_LIST} ${TARGET_NAME} + CACHE INTERNAL ROCM_CODEOBJ_LIST) + + # + # Build clang arguments into a string and tokenize it into a list + # The command "separate_arguments" will replace each instance of + # space char with a semi-colon char. Like any other program clang + # needs its arguments to be passed in as a list of tokens. The + # following strings are used to generate a code object and code + # asm files + # + string(CONCAT CODE_ARG_STR "-Xclang -finclude-default-header " + "-target amdgcn-amdh-amdhsa -mcpu=${TARGET_DEV} " + "${COMMON_BITCODE_LIB1} ${COMMON_BITCODE_LIB2} " + "${COMMON_BITCODE_LIB3} -cl-std=CL${OPENCL_VER} " + "${PROJECT_SOURCE_DIR}/${CL_FILE} -o ${KERNEL_DIR}/${CODEOBJ_FILE}") + string(CONCAT ASM_ARG_STR "-S -Xclang -finclude-default-header " + "-target amdgcn-amdh-amdhsa -mcpu=${TARGET_DEV} " + "${COMMON_BITCODE_LIB1} ${COMMON_BITCODE_LIB2} " + "${COMMON_BITCODE_LIB3} -cl-std=CL${OPENCL_VER} " + "${PROJECT_SOURCE_DIR}/${CL_FILE} -o ${KERNEL_DIR}/${CODEASM_FILE}") + set(ASM_ARG_LIST ${ASM_ARG_STR}) + set(CODE_ARG_LIST ${CODE_ARG_STR}) + separate_arguments(ASM_ARG_LIST) + separate_arguments(CODE_ARG_LIST) + + # + # Create a custom command to execute associated commands + # and a target it is associated with + # + add_custom_command(OUTPUT ${KERNEL_DIR}/${KNAME_EXE} + COMMAND ${CMAKE_COMMAND} -E make_directory ${KERNEL_DIR} + COMMAND ${CLANG} ${ASM_ARG_LIST} + COMMAND ${CLANG} ${CODE_ARG_LIST} + COMMENT "BUILDING KERNEL..." VERBATIM) + add_custom_target("${TARGET_NAME}" ALL DEPENDS "${KERNEL_DIR}/${KNAME_EXE}") + +endfunction(CompileKernel) + +function(buildCodeObjects kname) + + # Bind the name of CL file and associate + # a name for the target + set(KNAME_EXE "${kname}") + set(CL_FILE "${kname}_kernel.cl") + + # Iterate through list of target devices + foreach(tdev ${TARGET_DEV_LIST}) + CompileKernel(${kname} ${tdev}) + endforeach(tdev) + +endfunction(buildCodeObjects) + +buildCodeObjects("binary_search") + +# +# Create a custom target which will build the full suite +# of code objects for all kernels and target device pairs +# +add_custom_target(rocm_code_objs ALL DEPENDS ${ROCM_CODEOBJ_LIST}) + diff --git a/rocrtst/Kernels/binary_search_kernel.cl b/rocrtst/Kernels/binary_search_kernel.cl new file mode 100755 index 0000000000..eb3cca6c86 --- /dev/null +++ b/rocrtst/Kernels/binary_search_kernel.cl @@ -0,0 +1,127 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR + * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, + * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS WITH THE SOFTWARE. + * + */ + +/** + * One instance of this kernel call is a thread. + * Each thread finds out the segment in which it should look for the element. + * After that, it checks if the element is between the lower bound and upper + * bound of its segment. If yes, then this segment becomes the total + * searchspace for the next pass. + * + * To achieve this, it writes the lower bound and upper bound to the output + * array. In case the element at the left end (lower bound) matches the element + * we are looking for, that is marked in the output and we no longer need to + * look any further. + */ + +__kernel void +binarySearch(__global uint4 * outputArray, + __const __global uint2 * sortedArray, + const unsigned int findMe) { + unsigned int tid = get_global_id(0); + + // Then we find the elements for this thread + uint2 element = sortedArray[tid]; + + + // If the element to be found does not lie between + // them, then nothing left to do in this thread + if((element.x > findMe) || (element.y < findMe)) { + return; + } else { + // However, if the element does lie between the lower + // and upper bounds of this thread's searchspace + // we need to narrow down the search further in this + // search space + // The search space for this thread is marked in the + // output as being the total search space for the next pass + outputArray[0].x = tid; + outputArray[0].w = 1; + } +} + + +__kernel void +binarySearch_mulkeys(__global int *keys, + __global uint *input, + const unsigned int numKeys, + __global int *output) { + + int gid = get_global_id(0); + int lBound = gid * 256; + int uBound = lBound + 255; + + for(int i = 0; i < numKeys; i++) { + if(keys[i] >= input[lBound] && keys[i] <= input[uBound]) + output[i]=lBound; + } + +} + + +__kernel void +binarySearch_mulkeysConcurrent(__global uint *keys, + __global uint *input, + const unsigned int inputSize, // num. of inputs + const unsigned int numSubdivisions, + __global int *output) { + + int lBound = (get_global_id(0) % numSubdivisions) * (inputSize / numSubdivisions); + int uBound = lBound + inputSize / numSubdivisions; + int myKey = keys[get_global_id(0) / numSubdivisions]; + int mid; + + while(uBound >= lBound) { + mid = (lBound + uBound) / 2; + if(input[mid] == myKey) { + output[get_global_id(0) / numSubdivisions] = mid; + return; + } else if(input[mid] > myKey) { + uBound = mid - 1; + } else { + lBound = mid + 1; + } + } +}