AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Move Images code to hsa-runtime folder

Browse Source 
Change-Id: I53c1845d985ac3e9708d952865009c0021f3bb4f


[ROCm/ROCR-Runtime commit: 7e3db20826]
This commit is contained in:
Sean Keely
2020-04-30 00:18:36 -05:00
parent bd5ef0eff8
commit 1fc7f2dec7
70 changed files with 73713 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/CMakeLists.txt
+209
View File
@@ -0,0 +1,209 @@
#
# Minimum version of cmake required
#
cmake_minimum_required(VERSION 3.5.0)
#
# Required Defines on cmake command line
#
# 1) Set location of OpenCL header files
# OPENCL_DIR="Root for OpenCL install"
# If not set, the default value is "/opt/rocm/opencl"
#
# 2) Set location of CLANG/LLVM binary directory
# LLVM_DIR="Directory contains clang, llvm-link and llvm-dis
# If not set, the default value is "<PROJECT_BUILD_DIR>/lightning/bin"
#
# 3) Set BITCODE library directory
# BITCODE_DIR="Directory contains the bitcode library"
# If not set, the default value is "${OPENCL_DIR}/lib/x86_64/bitcode"
#
# 4) Set TARGET_DEVICES to indicate gpu types for kernel builds (e.g., "gfx803;gfx900; ...")
# If not set, the target devices are those have the Open Compute Library Controls (OCLC)
# bitcode file, "oclc_isa_version_*.amdgcn.bc", in the BITCODE directory
#
# Building - Should be automatic but for manual builds:
#
# 1) *** Create build folder e.g. "blit_src/build" - any name will do
# 2) Go to the build folder
# 3) Run "cmake .."
# 4) Run "make opencl_blit_objects.cpp"
#
## Include the cmake_modules utils.cmake
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../../../opensrc/hsa-runtime/cmake_modules" )
include ( utils )
# Flag to abort before executing after default initialization of cache variables
set (QUIT 0)
# Collect possible LLVM version directories.
set (LLVM_SEARCH_PATHS "")
set (LLVM_SEARCH_ROOT "${CMAKE_INSTALL_PREFIX}/llvm/lib/clang")
listsubdirs(${LLVM_SEARCH_ROOT} FOLDERS)
foreach(ITEM IN LISTS FOLDERS)
list (APPEND LLVM_SEARCH_PATHS "${LLVM_SEARCH_ROOT}/${ITEM}/include/")
endforeach()
if (NOT DEFINED OPENCL_VER)
set (OPENCL_VER "2.0")
endif()
set( OPENCL_VER ${OPENCL_VER} CACHE STRING "OpenCL version" FORCE )
get_include_path(BITCODE_DIR "Bitcode library path" RESULT FOUND NAMES "opencl.amdgcn.bc" HINTS "${CMAKE_INSTALL_PREFIX}/lib/bitcode" "${OPENCL_DIR}/lib/x86_64/bitcode")
if (NOT ${FOUND})
set (QUIT 1)
endif()
set (BITCODE_LIB "${BITCODE_DIR}/opencl.amdgcn.bc")
if (NOT EXISTS ${BITCODE_LIB})
message("ERROR: path to opencl.amdgcn.bc (${BITCODE_LIB}) is not valid. Is BITCODE_DIR correctly defined?")
set (QUIT 1)
endif()
get_include_path(LLVM_DIR "LLVM directory" RESULT FOUND NAMES "clang" HINTS "${CMAKE_INSTALL_PREFIX}/llvm/bin")
if (NOT ${FOUND})
set (QUIT 1)
endif()
set (CLANG "${LLVM_DIR}/clang")
if (NOT EXISTS ${CLANG})
message("ERROR: path to clang (${CLANG}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
set (LLVM_LINK "${LLVM_DIR}/llvm-link")
if (NOT EXISTS ${LLVM_LINK})
message("ERROR: path to llvm-link (${LLVM_LINK}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
set (LLVM_DIS "${LLVM_DIR}/llvm-dis")
if (NOT EXISTS ${LLVM_DIS})
message("ERROR: path to llvm-dis (${LLVM_DIS}) is not valid. Is LLVM_DIR correctly defined?")
set (QUIT 1)
endif()
# Value of Images Src Dir is bound in parent environment
set (KERNELS_DIR "${IMAGE_SOURCE_DIR}/blit_src")
# Define the target devices with xnack enable
if (NOT DEFINED XNACK_DEVS)
set (XNACK_DEVS "gfx801;gfx902")
endif()
set( XNACK_DEVS ${XNACK_DEVS} CACHE STRING "XNACK targets" FORCE )
# Determine the target devices if not specified
if (NOT DEFINED TARGET_DEVICES)
set (TARGET_DEVICES "gfx700;gfx701;gfx702;gfx801;gfx802;gfx803;gfx900;gfx902;gfx904;gfx906;gfx908;gfx1010;gfx1011;gfx1012")
endif()
set( TARGET_DEVICES ${TARGET_DEVICES} CACHE STRING "Build targets" FORCE )
# End of default configuration and path checking.
# Quit if configuration is incomplete.
if (QUIT)
message(FATAL_ERROR "Configuration halted.")
return()
endif()
# Determine the target triple
execute_process(COMMAND ${LLVM_DIS} ${BITCODE_LIB} -o - OUTPUT_VARIABLE LLVM_DIS_OUTPUT)
string(REGEX MATCH "(amdgcn-amd-.*)\"[\r\n]" QUOTED_TRIPLE "${LLVM_DIS_OUTPUT}")
string(REGEX REPLACE "[\"\r\n]" "" TARGET_TRIPLE "${QUOTED_TRIPLE}")
message("")
message("Build Setting:")
message(" Target Devices: ${TARGET_DEVICES}")
message(" Proj. Src Dir: ${PROJECT_SOURCE_DIR}")
message(" Proj. Bld Dir: ${PROJECT_BINARY_DIR}")
message(" Image Source Dir: ${IMAGE_SOURCE_DIR}")
message(" LLVM Dir: ${LLVM_DIR}")
message(" Clang path: ${CLANG}")
message(" OpenCL Dir: ${OPENCL_DIR}")
message(" OpenCL version: ${OPENCL_VER}")
message(" Bitcode Dir: ${BITCODE_DIR}")
message(" Target Triple: ${TARGET_TRIPLE}")
##==========================================
## Generate Kernel Bitcode
##==========================================
function(gen_kernel_bc TARGET_DEV XNACK_OPT FPREFIX INPUT_FILE OUTPUT_FILE)
string (REPLACE "gfx" "" GFXIP "${TARGET_DEV}")
separate_arguments(CLANG_ARG_LIST UNIX_COMMAND
"-O2 -x cl -target ${TARGET_TRIPLE} -Xclang -finclude-default-header -mcpu=${TARGET_DEV} -m${XNACK_OPT}
-nogpulib
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/opencl.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/ockl.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/ocml.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_daz_opt_on.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_isa_version_${GFXIP}.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_unsafe_math_off.amdgcn.bc
-Xclang -mlink-bitcode-file -Xclang ${BITCODE_DIR}/oclc_finite_only_off.amdgcn.bc
-cl-std=CL${OPENCL_VER} -o ${OUTPUT_FILE} ${INPUT_FILE}")
add_custom_target("${FPREFIX}" ${CLANG} ${CLANG_ARG_LIST}
COMMENT "BUILDING bitcode for ${FPREFIX}..."
VERBATIM)
message(" Kernel Source: " ${INPUT_FILE})
message(" Kernel Bitcode: " ${OUTPUT_FILE})
endfunction(gen_kernel_bc)
##==========================================
## Build the kernel for a device
##==========================================
function(build_kernel BLIT_NAME TARG_DEV)
list (FIND XNACK_DEVS ${TARG_DEV} XNACK_IDX)
if (${XNACK_IDX} GREATER -1)
set (XNACK_OPT "xnack")
else()
set (XNACK_OPT "no-xnack")
endif()
set (FILE_PREFIX "${BLIT_NAME}_${TARG_DEV}")
set (HSACO_TARG_LIST ${HSACO_TARG_LIST} "${FILE_PREFIX}" CACHE INTERNAL HSACO_TARG_LIST)
## generate kernel bitcodes
##
set (CL_FILE "${KERNELS_DIR}/imageblit_kernels.cl")
set (KERNEL_BC_FILE "${FILE_PREFIX}")
gen_kernel_bc(${TARG_DEV} ${XNACK_OPT} ${FILE_PREFIX} ${CL_FILE} ${KERNEL_BC_FILE})
endfunction(build_kernel)
##==========================================
## Build the kernel for a list of devices
##==========================================
function(build_kernel_for_devices BLIT_NAME)
set(HSACO_TARG_LIST PARENT_SCOPE)
foreach(dev ${TARGET_DEVICES})
message("\n Working on: ${dev} ...")
build_kernel(${BLIT_NAME} ${dev})
endforeach(dev)
endfunction(build_kernel_for_devices)
##==========================================
## Create BLIT Code Object blobs file
##==========================================
function(generate_blit_file BFILE)
file(REMOVE ${IMAGE_SOURCE_DIR}/${BFILE})
add_custom_command(OUTPUT ${IMAGE_SOURCE_DIR}/${BFILE}
COMMAND ${KERNELS_DIR}/create_hsaco_ascii_file.sh ${IMAGE_SOURCE_DIR}/${BFILE})
message("\n Will create ASCII bitcodes in ${BFILE} for ${TARGET_DEVICES} ... \n")
add_custom_target(${BFILE} DEPENDS ${HSACO_TARG_LIST} ${IMAGE_SOURCE_DIR}/${BFILE})
endfunction(generate_blit_file)
build_kernel_for_devices("ocl_blit_object")
generate_blit_file("opencl_blit_objects.cpp")
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/README.md
+61
View File
@@ -0,0 +1,61 @@
## OVERVIEW
This directory contains the CMakeLists.txt for automatically generating
the ASCII code object file, "opencl_blit_objects.cpp", which contains the
blobs of the code object of the Image BLIT kernels for the devices supported
on ROCm. The blobs are loaded by the image library and required to update
whenever a new device is introduced.
## ADD NEW DEVICE
To add a new supported device, the following steps are required:
1. Declare an extern variable of the device XXX, by adding the line of
"extern uint32_t ocl_blit_object_gfxNNN[];" in "blit_kernel.cpp"
2. Update the BlitKernel::GetPatchedBlitObject() function to support the
device by assigning "blit_code_object" to "ocl_blit_object_gfxNNN[]"
3. Add the gfxNNN to the TARGET_DEVICES list in CMakeLists.txt
4. If the new device requires XNACK, add it to the XNACK_DEVS list in CMakeLists.txt
5. Rebuild the image library
## REQUIREMENT
In order to create the code object file, the bitcodes of the kernels are
generated by the compiler and the following bitcode libraries are required,
opencl.amdgcn.bc
ocml.amdgcn.bc
irif.amdgcn.bc
oclc_correctly_rounded_sqrt_off.amdgcn.bc
oclc_daz_opt_on.amdgcn.bc
oclc_finite_only_off.amdgcn.bc
oclc_isa_version_<GFXIP>.amdgcn.bc
oclc_unsafe_math_off.amdgcn.bc
where <GFXIP> is the gfxip number of the GPU. The directory contains the
bitcode libraries is specified in a CMake varaible.
There are several variables are required for CMake to build the code
object file. All of them have default values, and defined as following:
OPENCL_DIR - the location of installed OpenCL
(Default: /opt/rocm/opencl)
BITCODE_DIR - the directory contains the bitcode library
(Default: ${OPENCL_DIR}/lib/x86_64/bitcode)
LLVM_DIR - the directory contains the clang, llvm-link and llvm-dis
executables
(Default: ${PROJECT_BUILD_DIR}/../lightning/bin)
TARGET_DEVICES - list of gpu types for kernel builds (eg. "gfx900;gfx902")
(Default: "gfx900;gfx902;gfx904")
## STEPS TO BUILD
$ make build
$ cd build
$ cmake -D${OPENCL_DIR} -D${BITCODE_DIR} -D${LLVM_DIR} -D${TARGET_DEVICES} ..
$ make opencl_blit_objects.cpp
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/create_hsaco_ascii_file.sh
Executable
+32
View File
@@ -0,0 +1,32 @@
#!/bin/bash -e
opencl_blit_file="$1"
if ! command -v xxd >/dev/null
then
echo "xxd not found!"
exit 1
fi
# Create the file in a temporary location and then move it in atomically
rm -rf "$opencl_blit_file.tmp"
{
cat <<EOF
//==============================================================================
// This file is automatically generated during build process, don't modify it
//==============================================================================
EOF
for file in ocl_blit_object*
do
xxd -i $file
echo -e '\n'
done
} > "$opencl_blit_file.tmp"
# Move the file atomically into place, so make doesn't get half a file
# but only if it has changed. cmp -s is happy for one file not to exist
cmp -s "$opencl_blit_file.tmp" "$opencl_blit_file" ||
mv -f "$opencl_blit_file.tmp" "$opencl_blit_file"
projects/rocr-runtime/runtime/hsa-runtime/image/blit_src/imageblit_kernels.cl
+615
View File
@@ -0,0 +1,615 @@
/// Kernel code for HSA image import/export/copy/clear in OpenCL C form.
uint4 read_image(__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
uint format,
int4 coords) {
switch (format) {
case 0: // 1D
return read_imageui(src1d, coords.x);
break;
case 1: // 2D
return read_imageui(src2d, coords.xy);
break;
case 2: // 3D
return read_imageui(src3d, coords);
break;
case 3: // 1DA
return read_imageui(src1da, coords.xy);
break;
case 4: // 2DA
return read_imageui(src2da, coords);
break;
// case 5: //1DB
// return read_imageui(src1db, coords.x);
// break;
default: // Critical failure.
return 0;
}
}
void write_image(__write_only image1d_t src1d,
__write_only image2d_t src2d,
__write_only image3d_t src3d,
__write_only image1d_array_t src1da,
__write_only image2d_array_t src2da,
uint format,
int4 coords,
uint4 texel) {
switch (format) {
case 0: // 1D
write_imageui(src1d, coords.x, texel);
break;
case 1: // 2D
write_imageui(src2d, coords.xy, texel);
break;
case 2: // 3D
write_imageui(src3d, coords, texel);
break;
case 3: // 1DA
write_imageui(src1da, coords.xy, texel);
break;
case 4: // 2DA
write_imageui(src2da, coords, texel);
break;
// case 5: //1DB
// write_imageui(src1db, coords.x, texel);
// break;
default: // Critical failure.
return;
}
}
float4 read_image_float(__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
uint format,
int4 coords) {
switch (format) {
case 0: // 1D
return read_imagef(src1d, coords.x);
break;
case 1: // 2D
return read_imagef(src2d, coords.xy);
break;
case 2: // 3D
return read_imagef(src3d, coords);
break;
case 3: // 1DA
return read_imagef(src1da, coords.xy);
break;
case 4: // 2DA
return read_imagef(src2da, coords);
break;
default: // Critical failure.
return 0;
}
}
void write_image_float(__write_only image1d_t src1d,
__write_only image2d_t src2d,
__write_only image3d_t src3d,
__write_only image1d_array_t src1da,
__write_only image2d_array_t src2da,
uint format,
int4 coords,
float4 texel) {
switch (format) {
case 0: // 1D
write_imagef(src1d, coords.x, texel);
break;
case 1: // 2D
write_imagef(src2d, coords.xy, texel);
break;
case 2: // 3D
write_imagef(src3d, coords, texel);
break;
case 3: // 1DA
write_imagef(src1da, coords.xy, texel);
break;
case 4: // 2DA
write_imagef(src2da, coords, texel);
break;
default: // Critical failure.
return;
}
}
void write_image_int(__write_only image1d_t src1d,
__write_only image2d_t src2d,
__write_only image3d_t src3d,
__write_only image1d_array_t src1da,
__write_only image2d_array_t src2da,
uint format,
int4 coords,
int4 texel) {
switch (format) {
case 0: // 1D
write_imagei(src1d, coords.x, texel);
break;
case 1: // 2D
write_imagei(src2d, coords.xy, texel);
break;
case 2: // 3D
write_imagei(src3d, coords, texel);
break;
case 3: // 1DA
write_imagei(src1da, coords.xy, texel);
break;
case 4: // 2DA
write_imagei(src2da, coords, texel);
break;
default: // Critical failure.
return;
}
}
//image handle is repeated since OCL doesn't allow pointers to or casting of images.
//dst is start of output pixel in destination buffer
//format.x is element count
//format.y is element size
//format.z is max(dword per pixel, 1)
//format.w is texture type.
//srcOrigin is start pixel address.
//No export for 64, 96, 128 bit formats
__kernel void copy_image_to_buffer(
__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
__global void* const dst,
int4 srcOrigin,
uint4 format,
ulong pitch,
ulong slice_pitch)
{
ulong idxDst;
int4 coordsSrc;
uint4 texel;
__global uchar* const dstUChar = (__global uchar* const)dst;
__global ushort* const dstUShort = (__global ushort* const)dst;
__global uint* const dstUInt = (__global uint* const)dst;
coordsSrc.x = get_global_id(0);
coordsSrc.y = get_global_id(1);
coordsSrc.z = get_global_id(2);
coordsSrc.w = 0;
idxDst = (coordsSrc.z * slice_pitch + coordsSrc.y * pitch +
coordsSrc.x) * format.z;
coordsSrc.x += srcOrigin.x;
coordsSrc.y += srcOrigin.y;
coordsSrc.z += srcOrigin.z;
texel = read_image(src1d, src2d, src3d, src1da, src2da, format.w, coordsSrc);
// Check components
switch (format.x) {
case 1:
// Check size
switch (format.y) {
case 1:
dstUChar[idxDst] = texel.x;
break;
case 2:
dstUShort[idxDst] = texel.x;
break;
case 4:
dstUInt[idxDst] = texel.x;
break;
}
break;
case 2:
// Check size
switch (format.y) {
case 1:
dstUShort[idxDst] = texel.x |
(texel.y << 8);
break;
case 2:
dstUInt[idxDst] = texel.x | (texel.y << 16);
break;
case 4:
dstUInt[idxDst++] = texel.x;
dstUInt[idxDst] = texel.y;
break;
}
break;
case 4:
// Check size
switch (format.y) {
case 1:
dstUInt[idxDst] = texel.x |
(texel.y << 8) |
(texel.z << 16) |
(texel.w << 24);
break;
case 2:
dstUInt[idxDst++] = texel.x | (texel.y << 16);
dstUInt[idxDst] = texel.z | (texel.w << 16);
break;
case 4:
dstUInt[idxDst++] = texel.x;
dstUInt[idxDst++] = texel.y;
dstUInt[idxDst++] = texel.z;
dstUInt[idxDst] = texel.w;
break;
}
break;
}
}
__kernel void copy_buffer_to_image(__global uint* src,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int4 dstOrigin,
uint4 format,
ulong pitch,
ulong slice_pitch) {
ulong idxSrc;
int4 coordsDst;
uint4 texel;
__global uint* srcUInt = src;
__global ushort* srcUShort = (__global ushort*)src;
__global uchar* srcUChar = (__global uchar*)src;
ushort tmpUShort;
uint tmpUInt;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
idxSrc = (coordsDst.z * slice_pitch + coordsDst.y * pitch + coordsDst.x) * format.z;
coordsDst.x += dstOrigin.x;
coordsDst.y += dstOrigin.y;
coordsDst.z += dstOrigin.z;
// Check components
switch (format.x) {
case 1:
// Check size
switch (format.y) {
case 1:
texel.x = (uint)srcUChar[idxSrc];
break;
case 2:
texel.x = (uint)srcUShort[idxSrc];
break;
case 4:
texel.x = srcUInt[idxSrc];
break;
}
break;
case 2:
// Check size
switch (format.y) {
case 1:
tmpUShort = srcUShort[idxSrc];
texel.x = (uint)(tmpUShort & 0xff);
texel.y = (uint)(tmpUShort >> 8);
break;
case 2:
tmpUInt = srcUInt[idxSrc];
texel.x = (tmpUInt & 0xffff);
texel.y = (tmpUInt >> 16);
break;
case 4:
texel.x = srcUInt[idxSrc++];
texel.y = srcUInt[idxSrc];
break;
}
break;
case 4:
// Check size
switch (format.y) {
case 1:
tmpUInt = srcUInt[idxSrc];
texel.x = tmpUInt & 0xff;
texel.y = (tmpUInt >> 8) & 0xff;
texel.z = (tmpUInt >> 16) & 0xff;
texel.w = (tmpUInt >> 24) & 0xff;
break;
case 2:
tmpUInt = srcUInt[idxSrc++];
texel.x = tmpUInt & 0xffff;
texel.y = (tmpUInt >> 16);
tmpUInt = srcUInt[idxSrc];
texel.z = tmpUInt & 0xffff;
texel.w = (tmpUInt >> 16);
break;
case 4:
texel.x = srcUInt[idxSrc++];
texel.y = srcUInt[idxSrc++];
texel.z = srcUInt[idxSrc++];
texel.w = srcUInt[idxSrc];
break;
}
break;
}
// Write the final pixel
write_image(dst1d, dst2d, dst3d, dst1da, dst2da, format.w, coordsDst, texel);
}
__kernel void copy_image_default(__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int4 srcOrigin,
int4 dstOrigin,
int srcFormat,
int dstFormat) {
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
uint4 texel;
texel = read_image(src1d, src2d, src3d, src1da, src2da, srcFormat, coordsSrc);
write_image(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coordsDst, texel);
}
float linear_to_standard_rgba(float l_val) {
float s_val = l_val;
if (isnan(s_val)) s_val = 0.0f;
if (s_val > 1.0f) {
s_val = 1.0f;
} else if (s_val < 0.0f) {
s_val = 0.0f;
} else if (s_val < 0.0031308f) {
s_val = 12.92f * s_val;
} else {
s_val = (1.055f * pow(s_val, 5.0f / 12.0f)) - 0.055f;
}
return s_val;
}
__kernel void copy_image_linear_to_standard(
__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
int srcFormat,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
int4 srcOrigin,
int4 dstOrigin) {
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
float4 texel;
texel = read_image_float(src1d, src2d, src3d, src1da, src2da, srcFormat, coordsSrc);
texel.x = linear_to_standard_rgba(texel.x);
texel.y = linear_to_standard_rgba(texel.y);
texel.z = linear_to_standard_rgba(texel.z);
write_image_float(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coordsDst, texel);
}
__kernel void copy_image_standard_to_linear(
__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
int srcFormat,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
int4 srcOrigin,
int4 dstOrigin) {
int4 coordsDst;
int4 coordsSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
coordsSrc = srcOrigin + coordsDst;
coordsDst += dstOrigin;
float4 texel;
texel = read_image_float(src1d, src2d, src3d, src1da, src2da, srcFormat, coordsSrc);
write_image_float(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coordsDst, texel);
}
__kernel void copy_image_1db(
__read_only image1d_buffer_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
int srcFormat,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
int4 srcOrigin,
int4 dstOrigin)
{
int coordDst;
int coordSrc;
coordDst = get_global_id(0);
coordSrc = srcOrigin.x + coordDst;
coordDst += dstOrigin.x;
uint4 texel;
texel = read_imageui(src1d, coordSrc);
write_imageui(dst1d, coordDst, texel);
}
__kernel void copy_image_1db_to_reg(
__read_only image1d_buffer_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
int srcFormat,
__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
int4 srcOrigin,
int4 dstOrigin)
{
int4 coordsDst;
int coordSrc;
coordsDst.x = get_global_id(0);
coordsDst.y = get_global_id(1);
coordsDst.z = get_global_id(2);
coordsDst.w = 0;
coordSrc = srcOrigin.x + coordsDst.x;
coordsDst += dstOrigin;
uint4 texel;
texel = read_imageui(src1d, coordSrc);
write_imageui(dst1d, coordsDst.x, texel);
}
__kernel void copy_image_reg_to_1db(
__read_only image1d_t src1d,
__read_only image2d_t src2d,
__read_only image3d_t src3d,
__read_only image1d_array_t src1da,
__read_only image2d_array_t src2da,
int srcFormat,
__write_only image1d_buffer_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
int4 srcOrigin,
int4 dstOrigin)
{
int coordDst;
int4 coordsSrc;
coordsSrc.x = get_global_id(0);
coordsSrc.y = get_global_id(1);
coordsSrc.z = get_global_id(2);
coordsSrc.w = 0;
coordDst = dstOrigin.x + coordsSrc.x;
coordsSrc += srcOrigin;
uint4 texel;
texel = read_imageui(src1d, coordsSrc.x);
write_imageui(dst1d, coordDst, texel);
}
__kernel void clear_image(__write_only image1d_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
uint type,
uint4 fill_data,
int4 origin) {
int4 coords;
coords.x = get_global_id(0);
coords.y = get_global_id(1);
coords.z = get_global_id(2);
coords.w = 0;
coords += origin;
// Check components
switch (type) {
case 0:
write_image_float(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coords, *(float4*)&fill_data);
break;
case 1:
write_image_int(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coords, *(int4*)&fill_data);
break;
case 2:
write_image(dst1d, dst2d, dst3d, dst1da, dst2da, dstFormat, coords, fill_data);
break;
}
}
__kernel void clear_image_1db(__write_only image1d_buffer_t dst1d,
__write_only image2d_t dst2d,
__write_only image3d_t dst3d,
__write_only image1d_array_t dst1da,
__write_only image2d_array_t dst2da,
int dstFormat,
uint4 fill_data,
int4 origin,
uint type) {
int4 coords;
coords.x = get_global_id(0);
coords += origin;
// Check components
switch (type) {
case 0:
write_imagef(dst1d, coords.x, *(float4*)&fill_data);
break;
case 1:
write_imagei(dst1d, coords.x, *(int4*)&fill_data);
break;
case 2:
write_imageui(dst1d, coords.x, fill_data);
break;
}
}