Native OpenCL C built-in functions. More...

#include "builtins.h"
#include "kernel.h"
#include "buffer.h"
#include "../events.h"
#include "../memobject.h"
#include <sys/mman.h>
#include <signal.h>
#include <llvm/Function.h>
#include <iostream>
#include <cstring>
#include <stdio.h>

Include dependency graph for builtins.cpp:

Go to the source code of this file.

Functions
unsigned char *	imageData (unsigned char *base, size_t x, size_t y, size_t z, size_t row_pitch, size_t slice_pitch, unsigned int bytes_per_pixel)
	Address of a pixel in an image.
void	setThreadLocalWorkGroup (Coal::CPUKernelWorkGroup *current)
	Set the current kernel work-group of this thread.
void *	getWorkItemsData (size_t &size)
	Work-item stacks.
void	setWorkItemsData (void *ptr, size_t size)
	Set work-item stacks.
static size_t	get_global_id (cl_uint dimindx)
static cl_uint	get_work_dim ()
static size_t	get_global_size (uint dimindx)
static size_t	get_local_size (uint dimindx)
static size_t	get_local_id (uint dimindx)
static size_t	get_num_groups (uint dimindx)
static size_t	get_group_id (uint dimindx)
static size_t	get_global_offset (uint dimindx)
static void	barrier (unsigned int flags)
static int	get_image_width (Image2D *image)
static int	get_image_height (Image2D *image)
static int	get_image_depth (Image3D *image)
static int	get_image_channel_data_type (Image2D *image)
static int	get_image_channel_order (Image2D *image)
static void *	image_data (Image2D image, int x, int y, int z, int order, int *type)
static bool	is_image_3d (Image3D *image)
static void	write_imagef (Image2D image, int x, int y, int z, float color)
static void	write_imagei (Image2D image, int x, int y, int z, int32_t color)
static void	write_imageui (Image2D image, int x, int y, int z, uint32_t color)
static void	read_imagefi (float result, Image2D image, int x, int y, int z, int32_t sampler)
static void	read_imageii (int32_t result, Image2D image, int x, int y, int z, int32_t sampler)
static void	read_imageuii (uint32_t result, Image2D image, int x, int y, int z, int32_t sampler)
static void	read_imageff (float result, Image2D image, float x, float y, float z, int32_t sampler)
static void	read_imageif (int32_t result, Image2D image, float x, float y, float z, int32_t sampler)
static void	read_imageuif (uint32_t result, Image2D image, float x, float y, float z, int32_t sampler)
static void	unimplemented_stub ()
void *	getBuiltin (const std::string &name)
	Return the address of a built-in function given its name.
Variables
__thread Coal::CPUKernelWorkGroup *	g_work_group
	`Coal::CPUKernelWorkGroup` currently running on this thread
__thread void *	work_items_data
	Space allocated for work-items stacks, see Implementing barriers.
__thread size_t	work_items_size
	Size of `work_items_data`, see Implementing barriers.

Detailed Description

Native OpenCL C built-in functions.

All these built-ins are directly called by kernels. When the LLVM JIT sees a function name it doesn't know, it calls getBuiltin() with this name as parameter. This function then returns the address of an actual function implementation, that finally gets called by the kernel when it is run.

Definition in file builtins.cpp.

Function Documentation

static void barrier ( unsigned int flags ) [static]

Definition at line 299 of file builtins.cpp.

References Coal::CPUKernelWorkGroup::barrier(), and g_work_group.

Referenced by getBuiltin().

static size_t get_global_id ( cl_uint dimindx ) [static]

Definition at line 259 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getGlobalId().

Referenced by getBuiltin().

static size_t get_global_offset ( uint dimindx ) [static]

Definition at line 294 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getGlobalOffset().

Referenced by getBuiltin().

static size_t get_global_size ( uint dimindx ) [static]

Definition at line 269 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getGlobalSize().

Referenced by getBuiltin().

static size_t get_group_id ( uint dimindx ) [static]

Definition at line 289 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getGroupID().

Referenced by getBuiltin().

static int get_image_channel_data_type ( Image2D * image ) [static]

Definition at line 324 of file builtins.cpp.

References Coal::Image2D::format().

Referenced by getBuiltin().

static int get_image_channel_order ( Image2D * image ) [static]

Definition at line 329 of file builtins.cpp.

References Coal::Image2D::format().

Referenced by getBuiltin().

static int get_image_depth ( Image3D * image ) [static]

Definition at line 316 of file builtins.cpp.

References Coal::Image3D::depth(), Coal::MemObject::Image3D, and Coal::Image3D::type().

Referenced by getBuiltin().

static int get_image_height ( Image2D * image ) [static]

Definition at line 311 of file builtins.cpp.

References Coal::Image2D::height().

Referenced by getBuiltin().

static int get_image_width ( Image2D * image ) [static]

Definition at line 306 of file builtins.cpp.

References Coal::Image2D::width().

Referenced by getBuiltin().

static size_t get_local_id ( uint dimindx ) [static]

Definition at line 279 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getLocalID().

Referenced by getBuiltin().

static size_t get_local_size ( uint dimindx ) [static]

Definition at line 274 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getLocalSize().

Referenced by getBuiltin().

static size_t get_num_groups ( uint dimindx ) [static]

Definition at line 284 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getNumGroups().

Referenced by getBuiltin().

static cl_uint get_work_dim ( ) [static]

Definition at line 264 of file builtins.cpp.

References g_work_group, and Coal::CPUKernelWorkGroup::getWorkDim().

Referenced by getBuiltin().

void* getBuiltin ( const std::string & name )

Return the address of a built-in function given its name.

Parameters:

name	name of the built-in whose address is requested

Definition at line 405 of file builtins.cpp.

References barrier(), Coal::CPUKernelWorkGroup::builtinNotFound(), g_work_group, get_global_id(), get_global_offset(), get_global_size(), get_group_id(), get_image_channel_data_type(), get_image_channel_order(), get_image_depth(), get_image_height(), get_image_width(), get_local_id(), get_local_size(), get_num_groups(), get_work_dim(), image_data(), is_image_3d(), read_imageff(), read_imagefi(), read_imageif(), read_imageii(), read_imageuif(), read_imageuii(), unimplemented_stub(), write_imagef(), write_imagei(), and write_imageui().

Referenced by Coal::CPUProgram::initJIT().

void* getWorkItemsData ( size_t & size )

Work-item stacks.

See also:: Implementing barriers

Parameters:

size	size of the allocated space for stacks

Returns:: address of the allocated space for stacks

Definition at line 82 of file builtins.cpp.

References work_items_data, and work_items_size.

Referenced by Coal::CPUKernelWorkGroup::barrier(), and worker().

static void* image_data	(	Image2D *	image,
		int	x,
		int	y,
		int	z,
		int *	order,
		int *	type
	)		`[static]`

Definition at line 334 of file builtins.cpp.

References Coal::Image2D::format(), g_work_group, and Coal::CPUKernelWorkGroup::getImageData().

Referenced by getBuiltin().

unsigned char* imageData	(	unsigned char *	base,
		size_t	x,
		size_t	y,
		size_t	z,
		size_t	row_pitch,
		size_t	slice_pitch,
		unsigned int	bytes_per_pixel
	)

Address of a pixel in an image.

This function is heavily used when Clover needs to address a pixel or a byte in a rectangular or three-dimensional image or buffer.

Parameters:

base	address of the first pixel in the image (address of the image itself)
x	X coordinate, cannot be bigger or equal to `width`
y	Y coordinate, cannot be bigger or equal to `height`
z	Z coordinate, cannot be bigger or equal to `depth` (1 for 2D arrays)
row_pitch	size in bytes of a row of pixels in the image
slice_pitch	size in bytes of a slice in a 3D array
bytes_per_pixel	bytes per pixel (1 for simple buffers), used when coordinates are in pixels and not in bytes.