Cmake project to build Rocm Code Objects
Change-Id: If3a631615316c203318bb5ae1df328a66e2919b1
Этот коммит содержится в:
Ramesh Errabolu
@@ -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})
|
||||
|
||||
Исполняемый файл
+127
@@ -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 <Name of Development Group, Name of Institution>,
|
||||
* 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;
|
||||
}
|
||||
}
|
||||
}
|
||||
Ссылка в новой задаче
Block a user