diff options
Diffstat (limited to 'src/gallium')
43 files changed, 8002 insertions, 1 deletions
diff --git a/src/gallium/state_trackers/Makefile b/src/gallium/state_trackers/Makefile index 0900efc664f..d5162c17507 100644 --- a/src/gallium/state_trackers/Makefile +++ b/src/gallium/state_trackers/Makefile @@ -17,7 +17,7 @@ subdirs: clean: - rm -f `find . -name \*.[oa]` + rm -f `find . -regex '.*\.l?[oa]'` rm -f `find . -name depend` diff --git a/src/gallium/state_trackers/clover/Doxyfile b/src/gallium/state_trackers/clover/Doxyfile new file mode 100644 index 00000000000..50250e75672 --- /dev/null +++ b/src/gallium/state_trackers/clover/Doxyfile @@ -0,0 +1,1716 @@ +# Doxyfile 1.7.4 + +# This file describes the settings to be used by the documentation system +# doxygen (www.doxygen.org) for a project. +# +# All text after a hash (#) is considered a comment and will be ignored. +# The format is: +# TAG = value [value, ...] +# For lists items can also be appended using: +# TAG += value [value, ...] +# Values that contain spaces should be placed between quotes (" "). + +#--------------------------------------------------------------------------- +# Project related configuration options +#--------------------------------------------------------------------------- + +# This tag specifies the encoding used for all characters in the config file +# that follow. The default is UTF-8 which is also the encoding used for all +# text before the first occurrence of this tag. Doxygen uses libiconv (or the +# iconv built into libc) for the transcoding. See +# http://www.gnu.org/software/libiconv for the list of possible encodings. + +DOXYFILE_ENCODING = UTF-8 + +# The PROJECT_NAME tag is a single word (or a sequence of words surrounded +# by quotes) that should identify the project. + +PROJECT_NAME = Clover + +# The PROJECT_NUMBER tag can be used to enter a project or revision number. +# This could be handy for archiving the generated documentation or +# if some version control system is used. + +PROJECT_NUMBER = + +# Using the PROJECT_BRIEF tag one can provide an optional one line description +# for a project that appears at the top of each page and should give viewer +# a quick idea about the purpose of the project. Keep the description short. + +PROJECT_BRIEF = + +# With the PROJECT_LOGO tag one can specify an logo or icon that is +# included in the documentation. The maximum height of the logo should not +# exceed 55 pixels and the maximum width should not exceed 200 pixels. +# Doxygen will copy the logo to the output directory. + +PROJECT_LOGO = + +# The OUTPUT_DIRECTORY tag is used to specify the (relative or absolute) +# base path where the generated documentation will be put. +# If a relative path is entered, it will be relative to the location +# where doxygen was started. If left blank the current directory will be used. + +OUTPUT_DIRECTORY = + +# If the CREATE_SUBDIRS tag is set to YES, then doxygen will create +# 4096 sub-directories (in 2 levels) under the output directory of each output +# format and will distribute the generated files over these directories. +# Enabling this option can be useful when feeding doxygen a huge amount of +# source files, where putting all generated files in the same directory would +# otherwise cause performance problems for the file system. + +CREATE_SUBDIRS = NO + +# The OUTPUT_LANGUAGE tag is used to specify the language in which all +# documentation generated by doxygen is written. Doxygen will use this +# information to generate all constant output in the proper language. +# The default language is English, other supported languages are: +# Afrikaans, Arabic, Brazilian, Catalan, Chinese, Chinese-Traditional, +# Croatian, Czech, Danish, Dutch, Esperanto, Farsi, Finnish, French, German, +# Greek, Hungarian, Italian, Japanese, Japanese-en (Japanese with English +# messages), Korean, Korean-en, Lithuanian, Norwegian, Macedonian, Persian, +# Polish, Portuguese, Romanian, Russian, Serbian, Serbian-Cyrillic, Slovak, +# Slovene, Spanish, Swedish, Ukrainian, and Vietnamese. + +OUTPUT_LANGUAGE = English + +# If the BRIEF_MEMBER_DESC tag is set to YES (the default) Doxygen will +# include brief member descriptions after the members that are listed in +# the file and class documentation (similar to JavaDoc). +# Set to NO to disable this. + +BRIEF_MEMBER_DESC = YES + +# If the REPEAT_BRIEF tag is set to YES (the default) Doxygen will prepend +# the brief description of a member or function before the detailed description. +# Note: if both HIDE_UNDOC_MEMBERS and BRIEF_MEMBER_DESC are set to NO, the +# brief descriptions will be completely suppressed. + +REPEAT_BRIEF = YES + +# This tag implements a quasi-intelligent brief description abbreviator +# that is used to form the text in various listings. Each string +# in this list, if found as the leading text of the brief description, will be +# stripped from the text and the result after processing the whole list, is +# used as the annotated text. Otherwise, the brief description is used as-is. +# If left blank, the following values are used ("$name" is automatically +# replaced with the name of the entity): "The $name class" "The $name widget" +# "The $name file" "is" "provides" "specifies" "contains" +# "represents" "a" "an" "the" + +ABBREVIATE_BRIEF = + +# If the ALWAYS_DETAILED_SEC and REPEAT_BRIEF tags are both set to YES then +# Doxygen will generate a detailed section even if there is only a brief +# description. + +ALWAYS_DETAILED_SEC = NO + +# If the INLINE_INHERITED_MEMB tag is set to YES, doxygen will show all +# inherited members of a class in the documentation of that class as if those +# members were ordinary class members. Constructors, destructors and assignment +# operators of the base classes will not be shown. + +INLINE_INHERITED_MEMB = NO + +# If the FULL_PATH_NAMES tag is set to YES then Doxygen will prepend the full +# path before files name in the file list and in the header files. If set +# to NO the shortest path that makes the file name unique will be used. + +FULL_PATH_NAMES = YES + +# If the FULL_PATH_NAMES tag is set to YES then the STRIP_FROM_PATH tag +# can be used to strip a user-defined part of the path. Stripping is +# only done if one of the specified strings matches the left-hand part of +# the path. The tag can be used to show relative paths in the file list. +# If left blank the directory from which doxygen is run is used as the +# path to strip. + +STRIP_FROM_PATH = + +# The STRIP_FROM_INC_PATH tag can be used to strip a user-defined part of +# the path mentioned in the documentation of a class, which tells +# the reader which header file to include in order to use a class. +# If left blank only the name of the header file containing the class +# definition is used. Otherwise one should specify the include paths that +# are normally passed to the compiler using the -I flag. + +STRIP_FROM_INC_PATH = + +# If the SHORT_NAMES tag is set to YES, doxygen will generate much shorter +# (but less readable) file names. This can be useful if your file system +# doesn't support long names like on DOS, Mac, or CD-ROM. + +SHORT_NAMES = NO + +# If the JAVADOC_AUTOBRIEF tag is set to YES then Doxygen +# will interpret the first line (until the first dot) of a JavaDoc-style +# comment as the brief description. If set to NO, the JavaDoc +# comments will behave just like regular Qt-style comments +# (thus requiring an explicit @brief command for a brief description.) + +JAVADOC_AUTOBRIEF = YES + +# If the QT_AUTOBRIEF tag is set to YES then Doxygen will +# interpret the first line (until the first dot) of a Qt-style +# comment as the brief description. If set to NO, the comments +# will behave just like regular Qt-style comments (thus requiring +# an explicit \brief command for a brief description.) + +QT_AUTOBRIEF = NO + +# The MULTILINE_CPP_IS_BRIEF tag can be set to YES to make Doxygen +# treat a multi-line C++ special comment block (i.e. a block of //! or /// +# comments) as a brief description. This used to be the default behaviour. +# The new default is to treat a multi-line C++ comment block as a detailed +# description. Set this tag to YES if you prefer the old behaviour instead. + +MULTILINE_CPP_IS_BRIEF = NO + +# If the INHERIT_DOCS tag is set to YES (the default) then an undocumented +# member inherits the documentation from any documented member that it +# re-implements. + +INHERIT_DOCS = YES + +# If the SEPARATE_MEMBER_PAGES tag is set to YES, then doxygen will produce +# a new page for each member. If set to NO, the documentation of a member will +# be part of the file/class/namespace that contains it. + +SEPARATE_MEMBER_PAGES = NO + +# The TAB_SIZE tag can be used to set the number of spaces in a tab. +# Doxygen uses this value to replace tabs by spaces in code fragments. + +TAB_SIZE = 8 + +# This tag can be used to specify a number of aliases that acts +# as commands in the documentation. An alias has the form "name=value". +# For example adding "sideeffect=\par Side Effects:\n" will allow you to +# put the command \sideeffect (or @sideeffect) in the documentation, which +# will result in a user-defined paragraph with heading "Side Effects:". +# You can put \n's in the value part of an alias to insert newlines. + +ALIASES = + +# Set the OPTIMIZE_OUTPUT_FOR_C tag to YES if your project consists of C +# sources only. Doxygen will then generate output that is more tailored for C. +# For instance, some of the names that are used will be different. The list +# of all members will be omitted, etc. + +OPTIMIZE_OUTPUT_FOR_C = NO + +# Set the OPTIMIZE_OUTPUT_JAVA tag to YES if your project consists of Java +# sources only. Doxygen will then generate output that is more tailored for +# Java. For instance, namespaces will be presented as packages, qualified +# scopes will look different, etc. + +OPTIMIZE_OUTPUT_JAVA = NO + +# Set the OPTIMIZE_FOR_FORTRAN tag to YES if your project consists of Fortran +# sources only. Doxygen will then generate output that is more tailored for +# Fortran. + +OPTIMIZE_FOR_FORTRAN = NO + +# Set the OPTIMIZE_OUTPUT_VHDL tag to YES if your project consists of VHDL +# sources. Doxygen will then generate output that is tailored for +# VHDL. + +OPTIMIZE_OUTPUT_VHDL = NO + +# Doxygen selects the parser to use depending on the extension of the files it +# parses. With this tag you can assign which parser to use for a given extension. +# Doxygen has a built-in mapping, but you can override or extend it using this +# tag. The format is ext=language, where ext is a file extension, and language +# is one of the parsers supported by doxygen: IDL, Java, Javascript, CSharp, C, +# C++, D, PHP, Objective-C, Python, Fortran, VHDL, C, C++. For instance to make +# doxygen treat .inc files as Fortran files (default is PHP), and .f files as C +# (default is Fortran), use: inc=Fortran f=C. Note that for custom extensions +# you also need to set FILE_PATTERNS otherwise the files are not read by doxygen. + +EXTENSION_MAPPING = + +# If you use STL classes (i.e. std::string, std::vector, etc.) but do not want +# to include (a tag file for) the STL sources as input, then you should +# set this tag to YES in order to let doxygen match functions declarations and +# definitions whose arguments contain STL classes (e.g. func(std::string); v.s. +# func(std::string) {}). This also makes the inheritance and collaboration +# diagrams that involve STL classes more complete and accurate. + +BUILTIN_STL_SUPPORT = YES + +# If you use Microsoft's C++/CLI language, you should set this option to YES to +# enable parsing support. + +CPP_CLI_SUPPORT = NO + +# Set the SIP_SUPPORT tag to YES if your project consists of sip sources only. +# Doxygen will parse them like normal C++ but will assume all classes use public +# instead of private inheritance when no explicit protection keyword is present. + +SIP_SUPPORT = NO + +# For Microsoft's IDL there are propget and propput attributes to indicate getter +# and setter methods for a property. Setting this option to YES (the default) +# will make doxygen replace the get and set methods by a property in the +# documentation. This will only work if the methods are indeed getting or +# setting a simple type. If this is not the case, or you want to show the +# methods anyway, you should set this option to NO. + +IDL_PROPERTY_SUPPORT = YES + +# If member grouping is used in the documentation and the DISTRIBUTE_GROUP_DOC +# tag is set to YES, then doxygen will reuse the documentation of the first +# member in the group (if any) for the other members of the group. By default +# all members of a group must be documented explicitly. + +DISTRIBUTE_GROUP_DOC = NO + +# Set the SUBGROUPING tag to YES (the default) to allow class member groups of +# the same type (for instance a group of public functions) to be put as a +# subgroup of that type (e.g. under the Public Functions section). Set it to +# NO to prevent subgrouping. Alternatively, this can be done per class using +# the \nosubgrouping command. + +SUBGROUPING = YES + +# When the INLINE_GROUPED_CLASSES tag is set to YES, classes, structs and +# unions are shown inside the group in which they are included (e.g. using +# @ingroup) instead of on a separate page (for HTML and Man pages) or +# section (for LaTeX and RTF). + +INLINE_GROUPED_CLASSES = NO + +# When TYPEDEF_HIDES_STRUCT is enabled, a typedef of a struct, union, or enum +# is documented as struct, union, or enum with the name of the typedef. So +# typedef struct TypeS {} TypeT, will appear in the documentation as a struct +# with name TypeT. When disabled the typedef will appear as a member of a file, +# namespace, or class. And the struct will be named TypeS. This can typically +# be useful for C code in case the coding convention dictates that all compound +# types are typedef'ed and only the typedef is referenced, never the tag name. + +TYPEDEF_HIDES_STRUCT = NO + +# The SYMBOL_CACHE_SIZE determines the size of the internal cache use to +# determine which symbols to keep in memory and which to flush to disk. +# When the cache is full, less often used symbols will be written to disk. +# For small to medium size projects (<1000 input files) the default value is +# probably good enough. For larger projects a too small cache size can cause +# doxygen to be busy swapping symbols to and from disk most of the time +# causing a significant performance penalty. +# If the system has enough physical memory increasing the cache will improve the +# performance by keeping more symbols in memory. Note that the value works on +# a logarithmic scale so increasing the size by one will roughly double the +# memory usage. The cache size is given by this formula: +# 2^(16+SYMBOL_CACHE_SIZE). The valid range is 0..9, the default is 0, +# corresponding to a cache size of 2^16 = 65536 symbols + +SYMBOL_CACHE_SIZE = 0 + +#--------------------------------------------------------------------------- +# Build related configuration options +#--------------------------------------------------------------------------- + +# If the EXTRACT_ALL tag is set to YES doxygen will assume all entities in +# documentation are documented, even if no documentation was available. +# Private class members and static file members will be hidden unless +# the EXTRACT_PRIVATE and EXTRACT_STATIC tags are set to YES + +EXTRACT_ALL = NO + +# If the EXTRACT_PRIVATE tag is set to YES all private members of a class +# will be included in the documentation. + +EXTRACT_PRIVATE = NO + +# If the EXTRACT_STATIC tag is set to YES all static members of a file +# will be included in the documentation. + +EXTRACT_STATIC = NO + +# If the EXTRACT_LOCAL_CLASSES tag is set to YES classes (and structs) +# defined locally in source files will be included in the documentation. +# If set to NO only classes defined in header files are included. + +EXTRACT_LOCAL_CLASSES = YES + +# This flag is only useful for Objective-C code. When set to YES local +# methods, which are defined in the implementation section but not in +# the interface are included in the documentation. +# If set to NO (the default) only methods in the interface are included. + +EXTRACT_LOCAL_METHODS = NO + +# If this flag is set to YES, the members of anonymous namespaces will be +# extracted and appear in the documentation as a namespace called +# 'anonymous_namespace{file}', where file will be replaced with the base +# name of the file that contains the anonymous namespace. By default +# anonymous namespaces are hidden. + +EXTRACT_ANON_NSPACES = YES + +# If the HIDE_UNDOC_MEMBERS tag is set to YES, Doxygen will hide all +# undocumented members of documented classes, files or namespaces. +# If set to NO (the default) these members will be included in the +# various overviews, but no documentation section is generated. +# This option has no effect if EXTRACT_ALL is enabled. + +HIDE_UNDOC_MEMBERS = NO + +# If the HIDE_UNDOC_CLASSES tag is set to YES, Doxygen will hide all +# undocumented classes that are normally visible in the class hierarchy. +# If set to NO (the default) these classes will be included in the various +# overviews. This option has no effect if EXTRACT_ALL is enabled. + +HIDE_UNDOC_CLASSES = NO + +# If the HIDE_FRIEND_COMPOUNDS tag is set to YES, Doxygen will hide all +# friend (class|struct|union) declarations. +# If set to NO (the default) these declarations will be included in the +# documentation. + +HIDE_FRIEND_COMPOUNDS = NO + +# If the HIDE_IN_BODY_DOCS tag is set to YES, Doxygen will hide any +# documentation blocks found inside the body of a function. +# If set to NO (the default) these blocks will be appended to the +# function's detailed documentation block. + +HIDE_IN_BODY_DOCS = NO + +# The INTERNAL_DOCS tag determines if documentation +# that is typed after a \internal command is included. If the tag is set +# to NO (the default) then the documentation will be excluded. +# Set it to YES to include the internal documentation. + +INTERNAL_DOCS = NO + +# If the CASE_SENSE_NAMES tag is set to NO then Doxygen will only generate +# file names in lower-case letters. If set to YES upper-case letters are also +# allowed. This is useful if you have classes or files whose names only differ +# in case and if your file system supports case sensitive file names. Windows +# and Mac users are advised to set this option to NO. + +CASE_SENSE_NAMES = YES + +# If the HIDE_SCOPE_NAMES tag is set to NO (the default) then Doxygen +# will show members with their full class and namespace scopes in the +# documentation. If set to YES the scope will be hidden. + +HIDE_SCOPE_NAMES = NO + +# If the SHOW_INCLUDE_FILES tag is set to YES (the default) then Doxygen +# will put a list of the files that are included by a file in the documentation +# of that file. + +SHOW_INCLUDE_FILES = YES + +# If the FORCE_LOCAL_INCLUDES tag is set to YES then Doxygen +# will list include files with double quotes in the documentation +# rather than with sharp brackets. + +FORCE_LOCAL_INCLUDES = NO + +# If the INLINE_INFO tag is set to YES (the default) then a tag [inline] +# is inserted in the documentation for inline members. + +INLINE_INFO = YES + +# If the SORT_MEMBER_DOCS tag is set to YES (the default) then doxygen +# will sort the (detailed) documentation of file and class members +# alphabetically by member name. If set to NO the members will appear in +# declaration order. + +SORT_MEMBER_DOCS = YES + +# If the SORT_BRIEF_DOCS tag is set to YES then doxygen will sort the +# brief documentation of file, namespace and class members alphabetically +# by member name. If set to NO (the default) the members will appear in +# declaration order. + +SORT_BRIEF_DOCS = NO + +# If the SORT_MEMBERS_CTORS_1ST tag is set to YES then doxygen +# will sort the (brief and detailed) documentation of class members so that +# constructors and destructors are listed first. If set to NO (the default) +# the constructors will appear in the respective orders defined by +# SORT_MEMBER_DOCS and SORT_BRIEF_DOCS. +# This tag will be ignored for brief docs if SORT_BRIEF_DOCS is set to NO +# and ignored for detailed docs if SORT_MEMBER_DOCS is set to NO. + +SORT_MEMBERS_CTORS_1ST = NO + +# If the SORT_GROUP_NAMES tag is set to YES then doxygen will sort the +# hierarchy of group names into alphabetical order. If set to NO (the default) +# the group names will appear in their defined order. + +SORT_GROUP_NAMES = NO + +# If the SORT_BY_SCOPE_NAME tag is set to YES, the class list will be +# sorted by fully-qualified names, including namespaces. If set to +# NO (the default), the class list will be sorted only by class name, +# not including the namespace part. +# Note: This option is not very useful if HIDE_SCOPE_NAMES is set to YES. +# Note: This option applies only to the class list, not to the +# alphabetical list. + +SORT_BY_SCOPE_NAME = NO + +# If the STRICT_PROTO_MATCHING option is enabled and doxygen fails to +# do proper type resolution of all parameters of a function it will reject a +# match between the prototype and the implementation of a member function even +# if there is only one candidate or it is obvious which candidate to choose +# by doing a simple string match. By disabling STRICT_PROTO_MATCHING doxygen +# will still accept a match between prototype and implementation in such cases. + +STRICT_PROTO_MATCHING = NO + +# The GENERATE_TODOLIST tag can be used to enable (YES) or +# disable (NO) the todo list. This list is created by putting \todo +# commands in the documentation. + +GENERATE_TODOLIST = YES + +# The GENERATE_TESTLIST tag can be used to enable (YES) or +# disable (NO) the test list. This list is created by putting \test +# commands in the documentation. + +GENERATE_TESTLIST = YES + +# The GENERATE_BUGLIST tag can be used to enable (YES) or +# disable (NO) the bug list. This list is created by putting \bug +# commands in the documentation. + +GENERATE_BUGLIST = YES + +# The GENERATE_DEPRECATEDLIST tag can be used to enable (YES) or +# disable (NO) the deprecated list. This list is created by putting +# \deprecated commands in the documentation. + +GENERATE_DEPRECATEDLIST= YES + +# The ENABLED_SECTIONS tag can be used to enable conditional +# documentation sections, marked by \if sectionname ... \endif. + +ENABLED_SECTIONS = + +# The MAX_INITIALIZER_LINES tag determines the maximum number of lines +# the initial value of a variable or macro consists of for it to appear in +# the documentation. If the initializer consists of more lines than specified +# here it will be hidden. Use a value of 0 to hide initializers completely. +# The appearance of the initializer of individual variables and macros in the +# documentation can be controlled using \showinitializer or \hideinitializer +# command in the documentation regardless of this setting. + +MAX_INITIALIZER_LINES = 30 + +# Set the SHOW_USED_FILES tag to NO to disable the list of files generated +# at the bottom of the documentation of classes and structs. If set to YES the +# list will mention the files that were used to generate the documentation. + +SHOW_USED_FILES = YES + +# If the sources in your project are distributed over multiple directories +# then setting the SHOW_DIRECTORIES tag to YES will show the directory hierarchy +# in the documentation. The default is NO. + +SHOW_DIRECTORIES = NO + +# Set the SHOW_FILES tag to NO to disable the generation of the Files page. +# This will remove the Files entry from the Quick Index and from the +# Folder Tree View (if specified). The default is YES. + +SHOW_FILES = YES + +# Set the SHOW_NAMESPACES tag to NO to disable the generation of the +# Namespaces page. +# This will remove the Namespaces entry from the Quick Index +# and from the Folder Tree View (if specified). The default is YES. + +SHOW_NAMESPACES = YES + +# The FILE_VERSION_FILTER tag can be used to specify a program or script that +# doxygen should invoke to get the current version for each file (typically from +# the version control system). Doxygen will invoke the program by executing (via +# popen()) the command <command> <input-file>, where <command> is the value of +# the FILE_VERSION_FILTER tag, and <input-file> is the name of an input file +# provided by doxygen. Whatever the program writes to standard output +# is used as the file version. See the manual for examples. + +FILE_VERSION_FILTER = + +# The LAYOUT_FILE tag can be used to specify a layout file which will be parsed +# by doxygen. The layout file controls the global structure of the generated +# output files in an output format independent way. The create the layout file +# that represents doxygen's defaults, run doxygen with the -l option. +# You can optionally specify a file name after the option, if omitted +# DoxygenLayout.xml will be used as the name of the layout file. + +LAYOUT_FILE = + +#--------------------------------------------------------------------------- +# configuration options related to warning and progress messages +#--------------------------------------------------------------------------- + +# The QUIET tag can be used to turn on/off the messages that are generated +# by doxygen. Possible values are YES and NO. If left blank NO is used. + +QUIET = NO + +# The WARNINGS tag can be used to turn on/off the warning messages that are +# generated by doxygen. Possible values are YES and NO. If left blank +# NO is used. + +WARNINGS = YES + +# If WARN_IF_UNDOCUMENTED is set to YES, then doxygen will generate warnings +# for undocumented members. If EXTRACT_ALL is set to YES then this flag will +# automatically be disabled. + +WARN_IF_UNDOCUMENTED = NO + +# If WARN_IF_DOC_ERROR is set to YES, doxygen will generate warnings for +# potential errors in the documentation, such as not documenting some +# parameters in a documented function, or documenting parameters that +# don't exist or using markup commands wrongly. + +WARN_IF_DOC_ERROR = YES + +# The WARN_NO_PARAMDOC option can be enabled to get warnings for +# functions that are documented, but have no documentation for their parameters +# or return value. If set to NO (the default) doxygen will only warn about +# wrong or incomplete parameter documentation, but not about the absence of +# documentation. + +WARN_NO_PARAMDOC = NO + +# The WARN_FORMAT tag determines the format of the warning messages that +# doxygen can produce. The string should contain the $file, $line, and $text +# tags, which will be replaced by the file and line number from which the +# warning originated and the warning text. Optionally the format may contain +# $version, which will be replaced by the version of the file (if it could +# be obtained via FILE_VERSION_FILTER) + +WARN_FORMAT = "$file:$line: $text" + +# The WARN_LOGFILE tag can be used to specify a file to which warning +# and error messages should be written. If left blank the output is written +# to stderr. + +WARN_LOGFILE = + +#--------------------------------------------------------------------------- +# configuration options related to the input files +#--------------------------------------------------------------------------- + +# The INPUT tag can be used to specify the files and/or directories that contain +# documented source files. You may enter file names like "myfile.cpp" or +# directories like "/usr/src/myproject". Separate the files or directories +# with spaces. + +INPUT = api/ core/ + +# This tag can be used to specify the character encoding of the source files +# that doxygen parses. Internally doxygen uses the UTF-8 encoding, which is +# also the default input encoding. Doxygen uses libiconv (or the iconv built +# into libc) for the transcoding. See http://www.gnu.org/software/libiconv for +# the list of possible encodings. + +INPUT_ENCODING = UTF-8 + +# If the value of the INPUT tag contains directories, you can use the +# FILE_PATTERNS tag to specify one or more wildcard pattern (like *.cpp +# and *.h) to filter out the source-files in the directories. If left +# blank the following patterns are tested: +# *.c *.cc *.cxx *.cpp *.c++ *.d *.java *.ii *.ixx *.ipp *.i++ *.inl *.h *.hh +# *.hxx *.hpp *.h++ *.idl *.odl *.cs *.php *.php3 *.inc *.m *.mm *.dox *.py +# *.f90 *.f *.for *.vhd *.vhdl + +FILE_PATTERNS = + +# The RECURSIVE tag can be used to turn specify whether or not subdirectories +# should be searched for input files as well. Possible values are YES and NO. +# If left blank NO is used. + +RECURSIVE = NO + +# The EXCLUDE tag can be used to specify files and/or directories that should +# excluded from the INPUT source files. This way you can easily exclude a +# subdirectory from a directory tree whose root is specified with the INPUT tag. + +EXCLUDE = + +# The EXCLUDE_SYMLINKS tag can be used select whether or not files or +# directories that are symbolic links (a Unix file system feature) are excluded +# from the input. + +EXCLUDE_SYMLINKS = NO + +# If the value of the INPUT tag contains directories, you can use the +# EXCLUDE_PATTERNS tag to specify one or more wildcard patterns to exclude +# certain files from those directories. Note that the wildcards are matched +# against the file with absolute path, so to exclude all test directories +# for example use the pattern */test/* + +EXCLUDE_PATTERNS = + +# The EXCLUDE_SYMBOLS tag can be used to specify one or more symbol names +# (namespaces, classes, functions, etc.) that should be excluded from the +# output. The symbol name can be a fully qualified name, a word, or if the +# wildcard * is used, a substring. Examples: ANamespace, AClass, +# AClass::ANamespace, ANamespace::*Test + +EXCLUDE_SYMBOLS = + +# The EXAMPLE_PATH tag can be used to specify one or more files or +# directories that contain example code fragments that are included (see +# the \include command). + +EXAMPLE_PATH = + +# If the value of the EXAMPLE_PATH tag contains directories, you can use the +# EXAMPLE_PATTERNS tag to specify one or more wildcard pattern (like *.cpp +# and *.h) to filter out the source-files in the directories. If left +# blank all files are included. + +EXAMPLE_PATTERNS = + +# If the EXAMPLE_RECURSIVE tag is set to YES then subdirectories will be +# searched for input files to be used with the \include or \dontinclude +# commands irrespective of the value of the RECURSIVE tag. +# Possible values are YES and NO. If left blank NO is used. + +EXAMPLE_RECURSIVE = NO + +# The IMAGE_PATH tag can be used to specify one or more files or +# directories that contain image that are included in the documentation (see +# the \image command). + +IMAGE_PATH = + +# The INPUT_FILTER tag can be used to specify a program that doxygen should +# invoke to filter for each input file. Doxygen will invoke the filter program +# by executing (via popen()) the command <filter> <input-file>, where <filter> +# is the value of the INPUT_FILTER tag, and <input-file> is the name of an +# input file. Doxygen will then use the output that the filter program writes +# to standard output. +# If FILTER_PATTERNS is specified, this tag will be +# ignored. + +INPUT_FILTER = + +# The FILTER_PATTERNS tag can be used to specify filters on a per file pattern +# basis. +# Doxygen will compare the file name with each pattern and apply the +# filter if there is a match. +# The filters are a list of the form: +# pattern=filter (like *.cpp=my_cpp_filter). See INPUT_FILTER for further +# info on how filters are used. If FILTER_PATTERNS is empty or if +# non of the patterns match the file name, INPUT_FILTER is applied. + +FILTER_PATTERNS = + +# If the FILTER_SOURCE_FILES tag is set to YES, the input filter (if set using +# INPUT_FILTER) will be used to filter the input files when producing source +# files to browse (i.e. when SOURCE_BROWSER is set to YES). + +FILTER_SOURCE_FILES = NO + +# The FILTER_SOURCE_PATTERNS tag can be used to specify source filters per file +# pattern. A pattern will override the setting for FILTER_PATTERN (if any) +# and it is also possible to disable source filtering for a specific pattern +# using *.ext= (so without naming a filter). This option only has effect when +# FILTER_SOURCE_FILES is enabled. + +FILTER_SOURCE_PATTERNS = + +#--------------------------------------------------------------------------- +# configuration options related to source browsing +#--------------------------------------------------------------------------- + +# If the SOURCE_BROWSER tag is set to YES then a list of source files will +# be generated. Documented entities will be cross-referenced with these sources. +# Note: To get rid of all source code in the generated output, make sure also +# VERBATIM_HEADERS is set to NO. + +SOURCE_BROWSER = NO + +# Setting the INLINE_SOURCES tag to YES will include the body +# of functions and classes directly in the documentation. + +INLINE_SOURCES = NO + +# Setting the STRIP_CODE_COMMENTS tag to YES (the default) will instruct +# doxygen to hide any special comment blocks from generated source code +# fragments. Normal C and C++ comments will always remain visible. + +STRIP_CODE_COMMENTS = YES + +# If the REFERENCED_BY_RELATION tag is set to YES +# then for each documented function all documented +# functions referencing it will be listed. + +REFERENCED_BY_RELATION = NO + +# If the REFERENCES_RELATION tag is set to YES +# then for each documented function all documented entities +# called/used by that function will be listed. + +REFERENCES_RELATION = NO + +# If the REFERENCES_LINK_SOURCE tag is set to YES (the default) +# and SOURCE_BROWSER tag is set to YES, then the hyperlinks from +# functions in REFERENCES_RELATION and REFERENCED_BY_RELATION lists will +# link to the source code. +# Otherwise they will link to the documentation. + +REFERENCES_LINK_SOURCE = YES + +# If the USE_HTAGS tag is set to YES then the references to source code +# will point to the HTML generated by the htags(1) tool instead of doxygen +# built-in source browser. The htags tool is part of GNU's global source +# tagging system (see http://www.gnu.org/software/global/global.html). You +# will need version 4.8.6 or higher. + +USE_HTAGS = NO + +# If the VERBATIM_HEADERS tag is set to YES (the default) then Doxygen +# will generate a verbatim copy of the header file for each class for +# which an include is specified. Set to NO to disable this. + +VERBATIM_HEADERS = YES + +#--------------------------------------------------------------------------- +# configuration options related to the alphabetical class index +#--------------------------------------------------------------------------- + +# If the ALPHABETICAL_INDEX tag is set to YES, an alphabetical index +# of all compounds will be generated. Enable this if the project +# contains a lot of classes, structs, unions or interfaces. + +ALPHABETICAL_INDEX = YES + +# If the alphabetical index is enabled (see ALPHABETICAL_INDEX) then +# the COLS_IN_ALPHA_INDEX tag can be used to specify the number of columns +# in which this list will be split (can be a number in the range [1..20]) + +COLS_IN_ALPHA_INDEX = 5 + +# In case all classes in a project start with a common prefix, all +# classes will be put under the same header in the alphabetical index. +# The IGNORE_PREFIX tag can be used to specify one or more prefixes that +# should be ignored while generating the index headers. + +IGNORE_PREFIX = + +#--------------------------------------------------------------------------- +# configuration options related to the HTML output +#--------------------------------------------------------------------------- + +# If the GENERATE_HTML tag is set to YES (the default) Doxygen will +# generate HTML output. + +GENERATE_HTML = YES + +# The HTML_OUTPUT tag is used to specify where the HTML docs will be put. +# If a relative path is entered the value of OUTPUT_DIRECTORY will be +# put in front of it. If left blank `html' will be used as the default path. + +HTML_OUTPUT = html + +# The HTML_FILE_EXTENSION tag can be used to specify the file extension for +# each generated HTML page (for example: .htm,.php,.asp). If it is left blank +# doxygen will generate files with .html extension. + +HTML_FILE_EXTENSION = .html + +# The HTML_HEADER tag can be used to specify a personal HTML header for +# each generated HTML page. If it is left blank doxygen will generate a +# standard header. Note that when using a custom header you are responsible +# for the proper inclusion of any scripts and style sheets that doxygen +# needs, which is dependent on the configuration options used. +# It is adviced to generate a default header using "doxygen -w html +# header.html footer.html stylesheet.css YourConfigFile" and then modify +# that header. Note that the header is subject to change so you typically +# have to redo this when upgrading to a newer version of doxygen or when changing the value of configuration settings such as GENERATE_TREEVIEW! + +HTML_HEADER = + +# The HTML_FOOTER tag can be used to specify a personal HTML footer for +# each generated HTML page. If it is left blank doxygen will generate a +# standard footer. + +HTML_FOOTER = + +# The HTML_STYLESHEET tag can be used to specify a user-defined cascading +# style sheet that is used by each HTML page. It can be used to +# fine-tune the look of the HTML output. If the tag is left blank doxygen +# will generate a default style sheet. Note that doxygen will try to copy +# the style sheet file to the HTML output directory, so don't put your own +# stylesheet in the HTML output directory as well, or it will be erased! + +HTML_STYLESHEET = + +# The HTML_EXTRA_FILES tag can be used to specify one or more extra images or +# other source files which should be copied to the HTML output directory. Note +# that these files will be copied to the base HTML output directory. Use the +# $relpath$ marker in the HTML_HEADER and/or HTML_FOOTER files to load these +# files. In the HTML_STYLESHEET file, use the file name only. Also note that +# the files will be copied as-is; there are no commands or markers available. + +HTML_EXTRA_FILES = + +# The HTML_COLORSTYLE_HUE tag controls the color of the HTML output. +# Doxygen will adjust the colors in the stylesheet and background images +# according to this color. Hue is specified as an angle on a colorwheel, +# see http://en.wikipedia.org/wiki/Hue for more information. +# For instance the value 0 represents red, 60 is yellow, 120 is green, +# 180 is cyan, 240 is blue, 300 purple, and 360 is red again. +# The allowed range is 0 to 359. + +HTML_COLORSTYLE_HUE = 220 + +# The HTML_COLORSTYLE_SAT tag controls the purity (or saturation) of +# the colors in the HTML output. For a value of 0 the output will use +# grayscales only. A value of 255 will produce the most vivid colors. + +HTML_COLORSTYLE_SAT = 100 + +# The HTML_COLORSTYLE_GAMMA tag controls the gamma correction applied to +# the luminance component of the colors in the HTML output. Values below +# 100 gradually make the output lighter, whereas values above 100 make +# the output darker. The value divided by 100 is the actual gamma applied, +# so 80 represents a gamma of 0.8, The value 220 represents a gamma of 2.2, +# and 100 does not change the gamma. + +HTML_COLORSTYLE_GAMMA = 80 + +# If the HTML_TIMESTAMP tag is set to YES then the footer of each generated HTML +# page will contain the date and time when the page was generated. Setting +# this to NO can help when comparing the output of multiple runs. + +HTML_TIMESTAMP = YES + +# If the HTML_ALIGN_MEMBERS tag is set to YES, the members of classes, +# files or namespaces will be aligned in HTML using tables. If set to +# NO a bullet list will be used. + +HTML_ALIGN_MEMBERS = YES + +# If the HTML_DYNAMIC_SECTIONS tag is set to YES then the generated HTML +# documentation will contain sections that can be hidden and shown after the +# page has loaded. For this to work a browser that supports +# JavaScript and DHTML is required (for instance Mozilla 1.0+, Firefox +# Netscape 6.0+, Internet explorer 5.0+, Konqueror, or Safari). + +HTML_DYNAMIC_SECTIONS = NO + +# If the GENERATE_DOCSET tag is set to YES, additional index files +# will be generated that can be used as input for Apple's Xcode 3 +# integrated development environment, introduced with OSX 10.5 (Leopard). +# To create a documentation set, doxygen will generate a Makefile in the +# HTML output directory. Running make will produce the docset in that +# directory and running "make install" will install the docset in +# ~/Library/Developer/Shared/Documentation/DocSets so that Xcode will find +# it at startup. +# See http://developer.apple.com/tools/creatingdocsetswithdoxygen.html +# for more information. + +GENERATE_DOCSET = NO + +# When GENERATE_DOCSET tag is set to YES, this tag determines the name of the +# feed. A documentation feed provides an umbrella under which multiple +# documentation sets from a single provider (such as a company or product suite) +# can be grouped. + +DOCSET_FEEDNAME = "Doxygen generated docs" + +# When GENERATE_DOCSET tag is set to YES, this tag specifies a string that +# should uniquely identify the documentation set bundle. This should be a +# reverse domain-name style string, e.g. com.mycompany.MyDocSet. Doxygen +# will append .docset to the name. + +DOCSET_BUNDLE_ID = org.doxygen.Project + +# When GENERATE_PUBLISHER_ID tag specifies a string that should uniquely identify +# the documentation publisher. This should be a reverse domain-name style +# string, e.g. com.mycompany.MyDocSet.documentation. + +DOCSET_PUBLISHER_ID = org.doxygen.Publisher + +# The GENERATE_PUBLISHER_NAME tag identifies the documentation publisher. + +DOCSET_PUBLISHER_NAME = Publisher + +# If the GENERATE_HTMLHELP tag is set to YES, additional index files +# will be generated that can be used as input for tools like the +# Microsoft HTML help workshop to generate a compiled HTML help file (.chm) +# of the generated HTML documentation. + +GENERATE_HTMLHELP = NO + +# If the GENERATE_HTMLHELP tag is set to YES, the CHM_FILE tag can +# be used to specify the file name of the resulting .chm file. You +# can add a path in front of the file if the result should not be +# written to the html output directory. + +CHM_FILE = + +# If the GENERATE_HTMLHELP tag is set to YES, the HHC_LOCATION tag can +# be used to specify the location (absolute path including file name) of +# the HTML help compiler (hhc.exe). If non-empty doxygen will try to run +# the HTML help compiler on the generated index.hhp. + +HHC_LOCATION = + +# If the GENERATE_HTMLHELP tag is set to YES, the GENERATE_CHI flag +# controls if a separate .chi index file is generated (YES) or that +# it should be included in the master .chm file (NO). + +GENERATE_CHI = NO + +# If the GENERATE_HTMLHELP tag is set to YES, the CHM_INDEX_ENCODING +# is used to encode HtmlHelp index (hhk), content (hhc) and project file +# content. + +CHM_INDEX_ENCODING = + +# If the GENERATE_HTMLHELP tag is set to YES, the BINARY_TOC flag +# controls whether a binary table of contents is generated (YES) or a +# normal table of contents (NO) in the .chm file. + +BINARY_TOC = NO + +# The TOC_EXPAND flag can be set to YES to add extra items for group members +# to the contents of the HTML help documentation and to the tree view. + +TOC_EXPAND = NO + +# If the GENERATE_QHP tag is set to YES and both QHP_NAMESPACE and +# QHP_VIRTUAL_FOLDER are set, an additional index file will be generated +# that can be used as input for Qt's qhelpgenerator to generate a +# Qt Compressed Help (.qch) of the generated HTML documentation. + +GENERATE_QHP = NO + +# If the QHG_LOCATION tag is specified, the QCH_FILE tag can +# be used to specify the file name of the resulting .qch file. +# The path specified is relative to the HTML output folder. + +QCH_FILE = + +# The QHP_NAMESPACE tag specifies the namespace to use when generating +# Qt Help Project output. For more information please see +# http://doc.trolltech.com/qthelpproject.html#namespace + +QHP_NAMESPACE = org.doxygen.Project + +# The QHP_VIRTUAL_FOLDER tag specifies the namespace to use when generating +# Qt Help Project output. For more information please see +# http://doc.trolltech.com/qthelpproject.html#virtual-folders + +QHP_VIRTUAL_FOLDER = doc + +# If QHP_CUST_FILTER_NAME is set, it specifies the name of a custom filter to +# add. For more information please see +# http://doc.trolltech.com/qthelpproject.html#custom-filters + +QHP_CUST_FILTER_NAME = + +# The QHP_CUST_FILT_ATTRS tag specifies the list of the attributes of the +# custom filter to add. For more information please see +# <a href="http://doc.trolltech.com/qthelpproject.html#custom-filters"> +# Qt Help Project / Custom Filters</a>. + +QHP_CUST_FILTER_ATTRS = + +# The QHP_SECT_FILTER_ATTRS tag specifies the list of the attributes this +# project's +# filter section matches. +# <a href="http://doc.trolltech.com/qthelpproject.html#filter-attributes"> +# Qt Help Project / Filter Attributes</a>. + +QHP_SECT_FILTER_ATTRS = + +# If the GENERATE_QHP tag is set to YES, the QHG_LOCATION tag can +# be used to specify the location of Qt's qhelpgenerator. +# If non-empty doxygen will try to run qhelpgenerator on the generated +# .qhp file. + +QHG_LOCATION = + +# If the GENERATE_ECLIPSEHELP tag is set to YES, additional index files +# will be generated, which together with the HTML files, form an Eclipse help +# plugin. To install this plugin and make it available under the help contents +# menu in Eclipse, the contents of the directory containing the HTML and XML +# files needs to be copied into the plugins directory of eclipse. The name of +# the directory within the plugins directory should be the same as +# the ECLIPSE_DOC_ID value. After copying Eclipse needs to be restarted before +# the help appears. + +GENERATE_ECLIPSEHELP = NO + +# A unique identifier for the eclipse help plugin. When installing the plugin +# the directory name containing the HTML and XML files should also have +# this name. + +ECLIPSE_DOC_ID = org.doxygen.Project + +# The DISABLE_INDEX tag can be used to turn on/off the condensed index at +# top of each HTML page. The value NO (the default) enables the index and +# the value YES disables it. + +DISABLE_INDEX = NO + +# The ENUM_VALUES_PER_LINE tag can be used to set the number of enum values +# (range [0,1..20]) that doxygen will group on one line in the generated HTML +# documentation. Note that a value of 0 will completely suppress the enum +# values from appearing in the overview section. + +ENUM_VALUES_PER_LINE = 4 + +# The GENERATE_TREEVIEW tag is used to specify whether a tree-like index +# structure should be generated to display hierarchical information. +# If the tag value is set to YES, a side panel will be generated +# containing a tree-like index structure (just like the one that +# is generated for HTML Help). For this to work a browser that supports +# JavaScript, DHTML, CSS and frames is required (i.e. any modern browser). +# Windows users are probably better off using the HTML help feature. + +GENERATE_TREEVIEW = NO + +# By enabling USE_INLINE_TREES, doxygen will generate the Groups, Directories, +# and Class Hierarchy pages using a tree view instead of an ordered list. + +USE_INLINE_TREES = NO + +# If the treeview is enabled (see GENERATE_TREEVIEW) then this tag can be +# used to set the initial width (in pixels) of the frame in which the tree +# is shown. + +TREEVIEW_WIDTH = 250 + +# When the EXT_LINKS_IN_WINDOW option is set to YES doxygen will open +# links to external symbols imported via tag files in a separate window. + +EXT_LINKS_IN_WINDOW = NO + +# Use this tag to change the font size of Latex formulas included +# as images in the HTML documentation. The default is 10. Note that +# when you change the font size after a successful doxygen run you need +# to manually remove any form_*.png images from the HTML output directory +# to force them to be regenerated. + +FORMULA_FONTSIZE = 10 + +# Use the FORMULA_TRANPARENT tag to determine whether or not the images +# generated for formulas are transparent PNGs. Transparent PNGs are +# not supported properly for IE 6.0, but are supported on all modern browsers. +# Note that when changing this option you need to delete any form_*.png files +# in the HTML output before the changes have effect. + +FORMULA_TRANSPARENT = YES + +# Enable the USE_MATHJAX option to render LaTeX formulas using MathJax +# (see http://www.mathjax.org) which uses client side Javascript for the +# rendering instead of using prerendered bitmaps. Use this if you do not +# have LaTeX installed or if you want to formulas look prettier in the HTML +# output. When enabled you also need to install MathJax separately and +# configure the path to it using the MATHJAX_RELPATH option. + +USE_MATHJAX = NO + +# When MathJax is enabled you need to specify the location relative to the +# HTML output directory using the MATHJAX_RELPATH option. The destination +# directory should contain the MathJax.js script. For instance, if the mathjax +# directory is located at the same level as the HTML output directory, then +# MATHJAX_RELPATH should be ../mathjax. The default value points to the +# mathjax.org site, so you can quickly see the result without installing +# MathJax, but it is strongly recommended to install a local copy of MathJax +# before deployment. + +MATHJAX_RELPATH = http://www.mathjax.org/mathjax + +# When the SEARCHENGINE tag is enabled doxygen will generate a search box +# for the HTML output. The underlying search engine uses javascript +# and DHTML and should work on any modern browser. Note that when using +# HTML help (GENERATE_HTMLHELP), Qt help (GENERATE_QHP), or docsets +# (GENERATE_DOCSET) there is already a search function so this one should +# typically be disabled. For large projects the javascript based search engine +# can be slow, then enabling SERVER_BASED_SEARCH may provide a better solution. + +SEARCHENGINE = YES + +# When the SERVER_BASED_SEARCH tag is enabled the search engine will be +# implemented using a PHP enabled web server instead of at the web client +# using Javascript. Doxygen will generate the search PHP script and index +# file to put on the web server. The advantage of the server +# based approach is that it scales better to large projects and allows +# full text search. The disadvantages are that it is more difficult to setup +# and does not have live searching capabilities. + +SERVER_BASED_SEARCH = NO + +#--------------------------------------------------------------------------- +# configuration options related to the LaTeX output +#--------------------------------------------------------------------------- + +# If the GENERATE_LATEX tag is set to YES (the default) Doxygen will +# generate Latex output. + +GENERATE_LATEX = YES + +# The LATEX_OUTPUT tag is used to specify where the LaTeX docs will be put. +# If a relative path is entered the value of OUTPUT_DIRECTORY will be +# put in front of it. If left blank `latex' will be used as the default path. + +LATEX_OUTPUT = latex + +# The LATEX_CMD_NAME tag can be used to specify the LaTeX command name to be +# invoked. If left blank `latex' will be used as the default command name. +# Note that when enabling USE_PDFLATEX this option is only used for +# generating bitmaps for formulas in the HTML output, but not in the +# Makefile that is written to the output directory. + +LATEX_CMD_NAME = latex + +# The MAKEINDEX_CMD_NAME tag can be used to specify the command name to +# generate index for LaTeX. If left blank `makeindex' will be used as the +# default command name. + +MAKEINDEX_CMD_NAME = makeindex + +# If the COMPACT_LATEX tag is set to YES Doxygen generates more compact +# LaTeX documents. This may be useful for small projects and may help to +# save some trees in general. + +COMPACT_LATEX = NO + +# The PAPER_TYPE tag can be used to set the paper type that is used +# by the printer. Possible values are: a4, letter, legal and +# executive. If left blank a4wide will be used. + +PAPER_TYPE = a4 + +# The EXTRA_PACKAGES tag can be to specify one or more names of LaTeX +# packages that should be included in the LaTeX output. + +EXTRA_PACKAGES = + +# The LATEX_HEADER tag can be used to specify a personal LaTeX header for +# the generated latex document. The header should contain everything until +# the first chapter. If it is left blank doxygen will generate a +# standard header. Notice: only use this tag if you know what you are doing! + +LATEX_HEADER = + +# The LATEX_FOOTER tag can be used to specify a personal LaTeX footer for +# the generated latex document. The footer should contain everything after +# the last chapter. If it is left blank doxygen will generate a +# standard footer. Notice: only use this tag if you know what you are doing! + +LATEX_FOOTER = + +# If the PDF_HYPERLINKS tag is set to YES, the LaTeX that is generated +# is prepared for conversion to pdf (using ps2pdf). The pdf file will +# contain links (just like the HTML output) instead of page references +# This makes the output suitable for online browsing using a pdf viewer. + +PDF_HYPERLINKS = YES + +# If the USE_PDFLATEX tag is set to YES, pdflatex will be used instead of +# plain latex in the generated Makefile. Set this option to YES to get a +# higher quality PDF documentation. + +USE_PDFLATEX = YES + +# If the LATEX_BATCHMODE tag is set to YES, doxygen will add the \\batchmode. +# command to the generated LaTeX files. This will instruct LaTeX to keep +# running if errors occur, instead of asking the user for help. +# This option is also used when generating formulas in HTML. + +LATEX_BATCHMODE = NO + +# If LATEX_HIDE_INDICES is set to YES then doxygen will not +# include the index chapters (such as File Index, Compound Index, etc.) +# in the output. + +LATEX_HIDE_INDICES = NO + +# If LATEX_SOURCE_CODE is set to YES then doxygen will include +# source code with syntax highlighting in the LaTeX output. +# Note that which sources are shown also depends on other settings +# such as SOURCE_BROWSER. + +LATEX_SOURCE_CODE = NO + +#--------------------------------------------------------------------------- +# configuration options related to the RTF output +#--------------------------------------------------------------------------- + +# If the GENERATE_RTF tag is set to YES Doxygen will generate RTF output +# The RTF output is optimized for Word 97 and may not look very pretty with +# other RTF readers or editors. + +GENERATE_RTF = NO + +# The RTF_OUTPUT tag is used to specify where the RTF docs will be put. +# If a relative path is entered the value of OUTPUT_DIRECTORY will be +# put in front of it. If left blank `rtf' will be used as the default path. + +RTF_OUTPUT = rtf + +# If the COMPACT_RTF tag is set to YES Doxygen generates more compact +# RTF documents. This may be useful for small projects and may help to +# save some trees in general. + +COMPACT_RTF = NO + +# If the RTF_HYPERLINKS tag is set to YES, the RTF that is generated +# will contain hyperlink fields. The RTF file will +# contain links (just like the HTML output) instead of page references. +# This makes the output suitable for online browsing using WORD or other +# programs which support those fields. +# Note: wordpad (write) and others do not support links. + +RTF_HYPERLINKS = NO + +# Load stylesheet definitions from file. Syntax is similar to doxygen's +# config file, i.e. a series of assignments. You only have to provide +# replacements, missing definitions are set to their default value. + +RTF_STYLESHEET_FILE = + +# Set optional variables used in the generation of an rtf document. +# Syntax is similar to doxygen's config file. + +RTF_EXTENSIONS_FILE = + +#--------------------------------------------------------------------------- +# configuration options related to the man page output +#--------------------------------------------------------------------------- + +# If the GENERATE_MAN tag is set to YES (the default) Doxygen will +# generate man pages + +GENERATE_MAN = NO + +# The MAN_OUTPUT tag is used to specify where the man pages will be put. +# If a relative path is entered the value of OUTPUT_DIRECTORY will be +# put in front of it. If left blank `man' will be used as the default path. + +MAN_OUTPUT = man + +# The MAN_EXTENSION tag determines the extension that is added to +# the generated man pages (default is the subroutine's section .3) + +MAN_EXTENSION = .3 + +# If the MAN_LINKS tag is set to YES and Doxygen generates man output, +# then it will generate one additional man file for each entity +# documented in the real man page(s). These additional files +# only source the real man page, but without them the man command +# would be unable to find the correct page. The default is NO. + +MAN_LINKS = NO + +#--------------------------------------------------------------------------- +# configuration options related to the XML output +#--------------------------------------------------------------------------- + +# If the GENERATE_XML tag is set to YES Doxygen will +# generate an XML file that captures the structure of +# the code including all documentation. + +GENERATE_XML = NO + +# The XML_OUTPUT tag is used to specify where the XML pages will be put. +# If a relative path is entered the value of OUTPUT_DIRECTORY will be +# put in front of it. If left blank `xml' will be used as the default path. + +XML_OUTPUT = xml + +# The XML_SCHEMA tag can be used to specify an XML schema, +# which can be used by a validating XML parser to check the +# syntax of the XML files. + +XML_SCHEMA = + +# The XML_DTD tag can be used to specify an XML DTD, +# which can be used by a validating XML parser to check the +# syntax of the XML files. + +XML_DTD = + +# If the XML_PROGRAMLISTING tag is set to YES Doxygen will +# dump the program listings (including syntax highlighting +# and cross-referencing information) to the XML output. Note that +# enabling this will significantly increase the size of the XML output. + +XML_PROGRAMLISTING = YES + +#--------------------------------------------------------------------------- +# configuration options for the AutoGen Definitions output +#--------------------------------------------------------------------------- + +# If the GENERATE_AUTOGEN_DEF tag is set to YES Doxygen will +# generate an AutoGen Definitions (see autogen.sf.net) file +# that captures the structure of the code including all +# documentation. Note that this feature is still experimental +# and incomplete at the moment. + +GENERATE_AUTOGEN_DEF = NO + +#--------------------------------------------------------------------------- +# configuration options related to the Perl module output +#--------------------------------------------------------------------------- + +# If the GENERATE_PERLMOD tag is set to YES Doxygen will +# generate a Perl module file that captures the structure of +# the code including all documentation. Note that this +# feature is still experimental and incomplete at the +# moment. + +GENERATE_PERLMOD = NO + +# If the PERLMOD_LATEX tag is set to YES Doxygen will generate +# the necessary Makefile rules, Perl scripts and LaTeX code to be able +# to generate PDF and DVI output from the Perl module output. + +PERLMOD_LATEX = NO + +# If the PERLMOD_PRETTY tag is set to YES the Perl module output will be +# nicely formatted so it can be parsed by a human reader. +# This is useful +# if you want to understand what is going on. +# On the other hand, if this +# tag is set to NO the size of the Perl module output will be much smaller +# and Perl will parse it just the same. + +PERLMOD_PRETTY = YES + +# The names of the make variables in the generated doxyrules.make file +# are prefixed with the string contained in PERLMOD_MAKEVAR_PREFIX. +# This is useful so different doxyrules.make files included by the same +# Makefile don't overwrite each other's variables. + +PERLMOD_MAKEVAR_PREFIX = + +#--------------------------------------------------------------------------- +# Configuration options related to the preprocessor +#--------------------------------------------------------------------------- + +# If the ENABLE_PREPROCESSING tag is set to YES (the default) Doxygen will +# evaluate all C-preprocessor directives found in the sources and include +# files. + +ENABLE_PREPROCESSING = YES + +# If the MACRO_EXPANSION tag is set to YES Doxygen will expand all macro +# names in the source code. If set to NO (the default) only conditional +# compilation will be performed. Macro expansion can be done in a controlled +# way by setting EXPAND_ONLY_PREDEF to YES. + +MACRO_EXPANSION = NO + +# If the EXPAND_ONLY_PREDEF and MACRO_EXPANSION tags are both set to YES +# then the macro expansion is limited to the macros specified with the +# PREDEFINED and EXPAND_AS_DEFINED tags. + +EXPAND_ONLY_PREDEF = NO + +# If the SEARCH_INCLUDES tag is set to YES (the default) the includes files +# pointed to by INCLUDE_PATH will be searched when a #include is found. + +SEARCH_INCLUDES = YES + +# The INCLUDE_PATH tag can be used to specify one or more directories that +# contain include files that are not input files but should be processed by +# the preprocessor. + +INCLUDE_PATH = + +# You can use the INCLUDE_FILE_PATTERNS tag to specify one or more wildcard +# patterns (like *.h and *.hpp) to filter out the header-files in the +# directories. If left blank, the patterns specified with FILE_PATTERNS will +# be used. + +INCLUDE_FILE_PATTERNS = + +# The PREDEFINED tag can be used to specify one or more macro names that +# are defined before the preprocessor is started (similar to the -D option of +# gcc). The argument of the tag is a list of macros of the form: name +# or name=definition (no spaces). If the definition and the = are +# omitted =1 is assumed. To prevent a macro definition from being +# undefined via #undef or recursively expanded use the := operator +# instead of the = operator. + +PREDEFINED = + +# If the MACRO_EXPANSION and EXPAND_ONLY_PREDEF tags are set to YES then +# this tag can be used to specify a list of macro names that should be expanded. +# The macro definition that is found in the sources will be used. +# Use the PREDEFINED tag if you want to use a different macro definition that +# overrules the definition found in the source code. + +EXPAND_AS_DEFINED = + +# If the SKIP_FUNCTION_MACROS tag is set to YES (the default) then +# doxygen's preprocessor will remove all references to function-like macros +# that are alone on a line, have an all uppercase name, and do not end with a +# semicolon, because these will confuse the parser if not removed. + +SKIP_FUNCTION_MACROS = YES + +#--------------------------------------------------------------------------- +# Configuration::additions related to external references +#--------------------------------------------------------------------------- + +# The TAGFILES option can be used to specify one or more tagfiles. +# Optionally an initial location of the external documentation +# can be added for each tagfile. The format of a tag file without +# this location is as follows: +# +# TAGFILES = file1 file2 ... +# Adding location for the tag files is done as follows: +# +# TAGFILES = file1=loc1 "file2 = loc2" ... +# where "loc1" and "loc2" can be relative or absolute paths or +# URLs. If a location is present for each tag, the installdox tool +# does not have to be run to correct the links. +# Note that each tag file must have a unique name +# (where the name does NOT include the path) +# If a tag file is not located in the directory in which doxygen +# is run, you must also specify the path to the tagfile here. + +TAGFILES = + +# When a file name is specified after GENERATE_TAGFILE, doxygen will create +# a tag file that is based on the input files it reads. + +GENERATE_TAGFILE = + +# If the ALLEXTERNALS tag is set to YES all external classes will be listed +# in the class index. If set to NO only the inherited external classes +# will be listed. + +ALLEXTERNALS = NO + +# If the EXTERNAL_GROUPS tag is set to YES all external groups will be listed +# in the modules index. If set to NO, only the current project's groups will +# be listed. + +EXTERNAL_GROUPS = YES + +# The PERL_PATH should be the absolute path and name of the perl script +# interpreter (i.e. the result of `which perl'). + +PERL_PATH = /usr/bin/perl + +#--------------------------------------------------------------------------- +# Configuration options related to the dot tool +#--------------------------------------------------------------------------- + +# If the CLASS_DIAGRAMS tag is set to YES (the default) Doxygen will +# generate a inheritance diagram (in HTML, RTF and LaTeX) for classes with base +# or super classes. Setting the tag to NO turns the diagrams off. Note that +# this option also works with HAVE_DOT disabled, but it is recommended to +# install and use dot, since it yields more powerful graphs. + +CLASS_DIAGRAMS = YES + +# You can define message sequence charts within doxygen comments using the \msc +# command. Doxygen will then run the mscgen tool (see +# http://www.mcternan.me.uk/mscgen/) to produce the chart and insert it in the +# documentation. The MSCGEN_PATH tag allows you to specify the directory where +# the mscgen tool resides. If left empty the tool is assumed to be found in the +# default search path. + +MSCGEN_PATH = + +# If set to YES, the inheritance and collaboration graphs will hide +# inheritance and usage relations if the target is undocumented +# or is not a class. + +HIDE_UNDOC_RELATIONS = YES + +# If you set the HAVE_DOT tag to YES then doxygen will assume the dot tool is +# available from the path. This tool is part of Graphviz, a graph visualization +# toolkit from AT&T and Lucent Bell Labs. The other options in this section +# have no effect if this option is set to NO (the default) + +HAVE_DOT = NO + +# The DOT_NUM_THREADS specifies the number of dot invocations doxygen is +# allowed to run in parallel. When set to 0 (the default) doxygen will +# base this on the number of processors available in the system. You can set it +# explicitly to a value larger than 0 to get control over the balance +# between CPU load and processing speed. + +DOT_NUM_THREADS = 0 + +# By default doxygen will write a font called Helvetica to the output +# directory and reference it in all dot files that doxygen generates. +# When you want a differently looking font you can specify the font name +# using DOT_FONTNAME. You need to make sure dot is able to find the font, +# which can be done by putting it in a standard location or by setting the +# DOTFONTPATH environment variable or by setting DOT_FONTPATH to the directory +# containing the font. + +DOT_FONTNAME = Helvetica + +# The DOT_FONTSIZE tag can be used to set the size of the font of dot graphs. +# The default size is 10pt. + +DOT_FONTSIZE = 10 + +# By default doxygen will tell dot to use the output directory to look for the +# FreeSans.ttf font (which doxygen will put there itself). If you specify a +# different font using DOT_FONTNAME you can set the path where dot +# can find it using this tag. + +DOT_FONTPATH = + +# If the CLASS_GRAPH and HAVE_DOT tags are set to YES then doxygen +# will generate a graph for each documented class showing the direct and +# indirect inheritance relations. Setting this tag to YES will force the +# the CLASS_DIAGRAMS tag to NO. + +CLASS_GRAPH = YES + +# If the COLLABORATION_GRAPH and HAVE_DOT tags are set to YES then doxygen +# will generate a graph for each documented class showing the direct and +# indirect implementation dependencies (inheritance, containment, and +# class references variables) of the class with other documented classes. + +COLLABORATION_GRAPH = YES + +# If the GROUP_GRAPHS and HAVE_DOT tags are set to YES then doxygen +# will generate a graph for groups, showing the direct groups dependencies + +GROUP_GRAPHS = YES + +# If the UML_LOOK tag is set to YES doxygen will generate inheritance and +# collaboration diagrams in a style similar to the OMG's Unified Modeling +# Language. + +UML_LOOK = NO + +# If set to YES, the inheritance and collaboration graphs will show the +# relations between templates and their instances. + +TEMPLATE_RELATIONS = NO + +# If the ENABLE_PREPROCESSING, SEARCH_INCLUDES, INCLUDE_GRAPH, and HAVE_DOT +# tags are set to YES then doxygen will generate a graph for each documented +# file showing the direct and indirect include dependencies of the file with +# other documented files. + +INCLUDE_GRAPH = YES + +# If the ENABLE_PREPROCESSING, SEARCH_INCLUDES, INCLUDED_BY_GRAPH, and +# HAVE_DOT tags are set to YES then doxygen will generate a graph for each +# documented header file showing the documented files that directly or +# indirectly include this file. + +INCLUDED_BY_GRAPH = YES + +# If the CALL_GRAPH and HAVE_DOT options are set to YES then +# doxygen will generate a call dependency graph for every global function +# or class method. Note that enabling this option will significantly increase +# the time of a run. So in most cases it will be better to enable call graphs +# for selected functions only using the \callgraph command. + +CALL_GRAPH = NO + +# If the CALLER_GRAPH and HAVE_DOT tags are set to YES then +# doxygen will generate a caller dependency graph for every global function +# or class method. Note that enabling this option will significantly increase +# the time of a run. So in most cases it will be better to enable caller +# graphs for selected functions only using the \callergraph command. + +CALLER_GRAPH = NO + +# If the GRAPHICAL_HIERARCHY and HAVE_DOT tags are set to YES then doxygen +# will generate a graphical hierarchy of all classes instead of a textual one. + +GRAPHICAL_HIERARCHY = YES + +# If the DIRECTORY_GRAPH, SHOW_DIRECTORIES and HAVE_DOT tags are set to YES +# then doxygen will show the dependencies a directory has on other directories +# in a graphical way. The dependency relations are determined by the #include +# relations between the files in the directories. + +DIRECTORY_GRAPH = YES + +# The DOT_IMAGE_FORMAT tag can be used to set the image format of the images +# generated by dot. Possible values are svg, png, jpg, or gif. +# If left blank png will be used. + +DOT_IMAGE_FORMAT = png + +# The tag DOT_PATH can be used to specify the path where the dot tool can be +# found. If left blank, it is assumed the dot tool can be found in the path. + +DOT_PATH = + +# The DOTFILE_DIRS tag can be used to specify one or more directories that +# contain dot files that are included in the documentation (see the +# \dotfile command). + +DOTFILE_DIRS = + +# The MSCFILE_DIRS tag can be used to specify one or more directories that +# contain msc files that are included in the documentation (see the +# \mscfile command). + +MSCFILE_DIRS = + +# The DOT_GRAPH_MAX_NODES tag can be used to set the maximum number of +# nodes that will be shown in the graph. If the number of nodes in a graph +# becomes larger than this value, doxygen will truncate the graph, which is +# visualized by representing a node as a red box. Note that doxygen if the +# number of direct children of the root node in a graph is already larger than +# DOT_GRAPH_MAX_NODES then the graph will not be shown at all. Also note +# that the size of a graph can be further restricted by MAX_DOT_GRAPH_DEPTH. + +DOT_GRAPH_MAX_NODES = 50 + +# The MAX_DOT_GRAPH_DEPTH tag can be used to set the maximum depth of the +# graphs generated by dot. A depth value of 3 means that only nodes reachable +# from the root by following a path via at most 3 edges will be shown. Nodes +# that lay further from the root node will be omitted. Note that setting this +# option to 1 or 2 may greatly reduce the computation time needed for large +# code bases. Also note that the size of a graph can be further restricted by +# DOT_GRAPH_MAX_NODES. Using a depth of 0 means no depth restriction. + +MAX_DOT_GRAPH_DEPTH = 0 + +# Set the DOT_TRANSPARENT tag to YES to generate images with a transparent +# background. This is disabled by default, because dot on Windows does not +# seem to support this out of the box. Warning: Depending on the platform used, +# enabling this option may lead to badly anti-aliased labels on the edges of +# a graph (i.e. they become hard to read). + +DOT_TRANSPARENT = NO + +# Set the DOT_MULTI_TARGETS tag to YES allow dot to generate multiple output +# files in one run (i.e. multiple -o and -T options on the command line). This +# makes dot run faster, but since only newer versions of dot (>1.8.10) +# support this, this feature is disabled by default. + +DOT_MULTI_TARGETS = NO + +# If the GENERATE_LEGEND tag is set to YES (the default) Doxygen will +# generate a legend page explaining the meaning of the various boxes and +# arrows in the dot generated graphs. + +GENERATE_LEGEND = YES + +# If the DOT_CLEANUP tag is set to YES (the default) Doxygen will +# remove the intermediate dot files that are used to generate +# the various graphs. + +DOT_CLEANUP = YES diff --git a/src/gallium/state_trackers/clover/Makefile.am b/src/gallium/state_trackers/clover/Makefile.am new file mode 100644 index 00000000000..da9f3bb92da --- /dev/null +++ b/src/gallium/state_trackers/clover/Makefile.am @@ -0,0 +1,71 @@ +AUTOMAKE_OPTIONS = subdir-objects + +AM_CPPFLAGS = \ + $(GALLIUM_PIPE_LOADER_DEFINES) \ + -DMESA_VERSION=\"$(MESA_VERSION)\" \ + -DPIPE_SEARCH_DIR=\"$(OPENCL_LIB_INSTALL_DIR)\" \ + -I$(top_srcdir)/include \ + -I$(top_srcdir)/src/gallium/include \ + -I$(top_srcdir)/src/gallium/drivers \ + -I$(top_srcdir)/src/gallium/auxiliary \ + -I$(top_srcdir)/src/gallium/winsys \ + -I$(srcdir) + +noinst_LTLIBRARIES = libclover.la libcltgsi.la libclllvm.la + +libcltgsi_la_CXXFLAGS = \ + -std=c++0x + +libcltgsi_la_SOURCES = \ + tgsi/compiler.cpp + +libclllvm_la_CXXFLAGS = \ + -std=c++98 + +libclllvm_la_SOURCES = \ + llvm/invocation.cpp + +libclover_la_CXXFLAGS = \ + -std=c++0x + +libclover_la_LIBADD = \ + libcltgsi.la libclllvm.la + +libclover_la_SOURCES = \ + core/base.hpp \ + core/compat.hpp \ + core/compiler.hpp \ + core/geometry.hpp \ + core/device.hpp \ + core/device.cpp \ + core/context.hpp \ + core/context.cpp \ + core/queue.hpp \ + core/queue.cpp \ + core/format.hpp \ + core/format.cpp \ + core/memory.hpp \ + core/memory.cpp \ + core/resource.hpp \ + core/resource.cpp \ + core/sampler.hpp \ + core/sampler.cpp \ + core/event.hpp \ + core/event.cpp \ + core/program.hpp \ + core/program.cpp \ + core/kernel.hpp \ + core/kernel.cpp \ + core/module.hpp \ + core/module.cpp \ + api/util.hpp \ + api/platform.cpp \ + api/device.cpp \ + api/context.cpp \ + api/queue.cpp \ + api/memory.cpp \ + api/transfer.cpp \ + api/sampler.cpp \ + api/event.cpp \ + api/program.cpp \ + api/kernel.cpp diff --git a/src/gallium/state_trackers/clover/api/context.cpp b/src/gallium/state_trackers/clover/api/context.cpp new file mode 100644 index 00000000000..c8d668933e5 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/context.cpp @@ -0,0 +1,120 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/context.hpp" + +using namespace clover; + +PUBLIC cl_context +clCreateContext(const cl_context_properties *props, cl_uint num_devs, + const cl_device_id *devs, + void (CL_CALLBACK *pfn_notify)(const char *, const void *, + size_t, void *), + void *user_data, cl_int *errcode_ret) try { + auto mprops = property_map(props); + + if (!devs || !num_devs || + (!pfn_notify && user_data)) + throw error(CL_INVALID_VALUE); + + if (any_of(is_zero<cl_device_id>(), devs, devs + num_devs)) + throw error(CL_INVALID_DEVICE); + + for (auto p : mprops) { + if (!(p.first == CL_CONTEXT_PLATFORM && + (cl_platform_id)p.second == NULL)) + throw error(CL_INVALID_PROPERTY); + } + + ret_error(errcode_ret, CL_SUCCESS); + return new context( + property_vector(mprops), + std::vector<cl_device_id>(devs, devs + num_devs)); + +} catch(error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_context +clCreateContextFromType(const cl_context_properties *props, + cl_device_type type, + void (CL_CALLBACK *pfn_notify)( + const char *, const void *, size_t, void *), + void *user_data, cl_int *errcode_ret) { + cl_device_id dev; + cl_int ret; + + ret = clGetDeviceIDs(0, type, 1, &dev, 0); + if (ret) { + ret_error(errcode_ret, ret); + return NULL; + } + + return clCreateContext(props, 1, &dev, pfn_notify, user_data, errcode_ret); +} + +PUBLIC cl_int +clRetainContext(cl_context ctx) { + if (!ctx) + return CL_INVALID_CONTEXT; + + ctx->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseContext(cl_context ctx) { + if (!ctx) + return CL_INVALID_CONTEXT; + + if (ctx->release()) + delete ctx; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetContextInfo(cl_context ctx, cl_context_info param, + size_t size, void *buf, size_t *size_ret) { + if (!ctx) + return CL_INVALID_CONTEXT; + + switch (param) { + case CL_CONTEXT_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, ctx->ref_count()); + + case CL_CONTEXT_NUM_DEVICES: + return scalar_property<cl_uint>(buf, size, size_ret, ctx->devs.size()); + + case CL_CONTEXT_DEVICES: + return vector_property<cl_device_id>(buf, size, size_ret, ctx->devs); + + case CL_CONTEXT_PROPERTIES: + return vector_property<cl_context_properties>(buf, size, size_ret, + ctx->props()); + + default: + return CL_INVALID_VALUE; + } +} diff --git a/src/gallium/state_trackers/clover/api/device.cpp b/src/gallium/state_trackers/clover/api/device.cpp new file mode 100644 index 00000000000..03767519aaf --- /dev/null +++ b/src/gallium/state_trackers/clover/api/device.cpp @@ -0,0 +1,262 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/device.hpp" + +using namespace clover; + +static device_registry registry; + +PUBLIC cl_int +clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type, + cl_uint num_entries, cl_device_id *devices, + cl_uint *num_devices) { + std::vector<cl_device_id> devs; + + if (platform != NULL) + return CL_INVALID_PLATFORM; + + if ((!num_entries && devices) || + (!num_devices && !devices)) + return CL_INVALID_VALUE; + + // Collect matching devices + for (device &dev : registry) { + if (((device_type & CL_DEVICE_TYPE_DEFAULT) && + &dev == ®istry.front()) || + (device_type & dev.type())) + devs.push_back(&dev); + } + + if (devs.empty()) + return CL_DEVICE_NOT_FOUND; + + // ...and return the requested data. + if (num_devices) + *num_devices = devs.size(); + if (devices) + std::copy_n(devs.begin(), + std::min((cl_uint)devs.size(), num_entries), + devices); + + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetDeviceInfo(cl_device_id dev, cl_device_info param, + size_t size, void *buf, size_t *size_ret) { + if (!dev) + return CL_INVALID_DEVICE; + + switch (param) { + case CL_DEVICE_TYPE: + return scalar_property<cl_device_type>(buf, size, size_ret, dev->type()); + + case CL_DEVICE_VENDOR_ID: + return scalar_property<cl_uint>(buf, size, size_ret, dev->vendor_id()); + + case CL_DEVICE_MAX_COMPUTE_UNITS: + return scalar_property<cl_uint>(buf, size, size_ret, 1); + + case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: + return scalar_property<cl_uint>(buf, size, size_ret, + dev->max_block_size().size()); + + case CL_DEVICE_MAX_WORK_ITEM_SIZES: + return vector_property<size_t>(buf, size, size_ret, + dev->max_block_size()); + + case CL_DEVICE_MAX_WORK_GROUP_SIZE: + return scalar_property<size_t>(buf, size, size_ret, SIZE_MAX); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR: + return scalar_property<cl_uint>(buf, size, size_ret, 16); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT: + return scalar_property<cl_uint>(buf, size, size_ret, 8); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT: + return scalar_property<cl_uint>(buf, size, size_ret, 4); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG: + return scalar_property<cl_uint>(buf, size, size_ret, 2); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT: + return scalar_property<cl_uint>(buf, size, size_ret, 4); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE: + return scalar_property<cl_uint>(buf, size, size_ret, 2); + + case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF: + return scalar_property<cl_uint>(buf, size, size_ret, 0); + + case CL_DEVICE_MAX_CLOCK_FREQUENCY: + return scalar_property<cl_uint>(buf, size, size_ret, 0); + + case CL_DEVICE_ADDRESS_BITS: + return scalar_property<cl_uint>(buf, size, size_ret, 32); + + case CL_DEVICE_MAX_READ_IMAGE_ARGS: + return scalar_property<cl_uint>(buf, size, size_ret, + dev->max_images_read()); + + case CL_DEVICE_MAX_WRITE_IMAGE_ARGS: + return scalar_property<cl_uint>(buf, size, size_ret, + dev->max_images_write()); + + case CL_DEVICE_MAX_MEM_ALLOC_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, 0); + + case CL_DEVICE_IMAGE2D_MAX_WIDTH: + case CL_DEVICE_IMAGE2D_MAX_HEIGHT: + return scalar_property<size_t>(buf, size, size_ret, + 1 << dev->max_image_levels_2d()); + + case CL_DEVICE_IMAGE3D_MAX_WIDTH: + case CL_DEVICE_IMAGE3D_MAX_HEIGHT: + case CL_DEVICE_IMAGE3D_MAX_DEPTH: + return scalar_property<size_t>(buf, size, size_ret, + 1 << dev->max_image_levels_3d()); + + case CL_DEVICE_IMAGE_SUPPORT: + return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE); + + case CL_DEVICE_MAX_PARAMETER_SIZE: + return scalar_property<size_t>(buf, size, size_ret, + dev->max_mem_input()); + + case CL_DEVICE_MAX_SAMPLERS: + return scalar_property<cl_uint>(buf, size, size_ret, + dev->max_samplers()); + + case CL_DEVICE_MEM_BASE_ADDR_ALIGN: + case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE: + return scalar_property<cl_uint>(buf, size, size_ret, 128); + + case CL_DEVICE_SINGLE_FP_CONFIG: + return scalar_property<cl_device_fp_config>(buf, size, size_ret, + CL_FP_DENORM | CL_FP_INF_NAN | CL_FP_ROUND_TO_NEAREST); + + case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE: + return scalar_property<cl_device_mem_cache_type>(buf, size, size_ret, + CL_NONE); + + case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE: + return scalar_property<cl_uint>(buf, size, size_ret, 0); + + case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, 0); + + case CL_DEVICE_GLOBAL_MEM_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, + dev->max_mem_global()); + + case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, + dev->max_const_buffer_size()); + + case CL_DEVICE_MAX_CONSTANT_ARGS: + return scalar_property<cl_uint>(buf, size, size_ret, + dev->max_const_buffers()); + + case CL_DEVICE_LOCAL_MEM_TYPE: + return scalar_property<cl_device_local_mem_type>(buf, size, size_ret, + CL_LOCAL); + + case CL_DEVICE_LOCAL_MEM_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, + dev->max_mem_local()); + + case CL_DEVICE_ERROR_CORRECTION_SUPPORT: + return scalar_property<cl_bool>(buf, size, size_ret, CL_FALSE); + + case CL_DEVICE_PROFILING_TIMER_RESOLUTION: + return scalar_property<size_t>(buf, size, size_ret, 0); + + case CL_DEVICE_ENDIAN_LITTLE: + return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE); + + case CL_DEVICE_AVAILABLE: + case CL_DEVICE_COMPILER_AVAILABLE: + return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE); + + case CL_DEVICE_EXECUTION_CAPABILITIES: + return scalar_property<cl_device_exec_capabilities>(buf, size, size_ret, + CL_EXEC_KERNEL); + + case CL_DEVICE_QUEUE_PROPERTIES: + return scalar_property<cl_command_queue_properties>(buf, size, size_ret, + CL_QUEUE_PROFILING_ENABLE); + + case CL_DEVICE_NAME: + return string_property(buf, size, size_ret, dev->device_name()); + + case CL_DEVICE_VENDOR: + return string_property(buf, size, size_ret, dev->vendor_name()); + + case CL_DRIVER_VERSION: + return string_property(buf, size, size_ret, MESA_VERSION); + + case CL_DEVICE_PROFILE: + return string_property(buf, size, size_ret, "FULL_PROFILE"); + + case CL_DEVICE_VERSION: + return string_property(buf, size, size_ret, "OpenCL 1.1 MESA " MESA_VERSION); + + case CL_DEVICE_EXTENSIONS: + return string_property(buf, size, size_ret, ""); + + case CL_DEVICE_PLATFORM: + return scalar_property<cl_platform_id>(buf, size, size_ret, NULL); + + case CL_DEVICE_HOST_UNIFIED_MEMORY: + return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR: + return scalar_property<cl_uint>(buf, size, size_ret, 16); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT: + return scalar_property<cl_uint>(buf, size, size_ret, 8); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT: + return scalar_property<cl_uint>(buf, size, size_ret, 4); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG: + return scalar_property<cl_uint>(buf, size, size_ret, 2); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT: + return scalar_property<cl_uint>(buf, size, size_ret, 4); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE: + return scalar_property<cl_uint>(buf, size, size_ret, 2); + + case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF: + return scalar_property<cl_uint>(buf, size, size_ret, 0); + + case CL_DEVICE_OPENCL_C_VERSION: + return string_property(buf, size, size_ret, "OpenCL C 1.1"); + + default: + return CL_INVALID_VALUE; + } +} diff --git a/src/gallium/state_trackers/clover/api/event.cpp b/src/gallium/state_trackers/clover/api/event.cpp new file mode 100644 index 00000000000..d6c37f6aef2 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/event.cpp @@ -0,0 +1,239 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/event.hpp" + +using namespace clover; + +PUBLIC cl_event +clCreateUserEvent(cl_context ctx, cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + ret_error(errcode_ret, CL_SUCCESS); + return new soft_event(*ctx, {}, false); + +} catch(error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clSetUserEventStatus(cl_event ev, cl_int status) { + if (!dynamic_cast<soft_event *>(ev)) + return CL_INVALID_EVENT; + + if (status > 0) + return CL_INVALID_VALUE; + + if (ev->status() <= 0) + return CL_INVALID_OPERATION; + + if (status) + ev->abort(status); + else + ev->trigger(); + + return CL_SUCCESS; +} + +PUBLIC cl_int +clWaitForEvents(cl_uint num_evs, const cl_event *evs) try { + if (!num_evs || !evs) + throw error(CL_INVALID_VALUE); + + std::for_each(evs, evs + num_evs, [&](const cl_event ev) { + if (!ev) + throw error(CL_INVALID_EVENT); + + if (&ev->ctx != &evs[0]->ctx) + throw error(CL_INVALID_CONTEXT); + + if (ev->status() < 0) + throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST); + }); + + // Create a temporary soft event that depends on all the events in + // the wait list + ref_ptr<soft_event> sev = transfer( + new soft_event(evs[0]->ctx, { evs, evs + num_evs }, true)); + + // ...and wait on it. + sev->wait(); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clGetEventInfo(cl_event ev, cl_event_info param, + size_t size, void *buf, size_t *size_ret) { + if (!ev) + return CL_INVALID_EVENT; + + switch (param) { + case CL_EVENT_COMMAND_QUEUE: + return scalar_property<cl_command_queue>(buf, size, size_ret, ev->queue()); + + case CL_EVENT_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, &ev->ctx); + + case CL_EVENT_COMMAND_TYPE: + return scalar_property<cl_command_type>(buf, size, size_ret, ev->command()); + + case CL_EVENT_COMMAND_EXECUTION_STATUS: + return scalar_property<cl_int>(buf, size, size_ret, ev->status()); + + case CL_EVENT_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, ev->ref_count()); + + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clSetEventCallback(cl_event ev, cl_int type, + void (CL_CALLBACK *pfn_event_notify)(cl_event, cl_int, + void *), + void *user_data) try { + if (!ev) + throw error(CL_INVALID_EVENT); + + if (!pfn_event_notify || type != CL_COMPLETE) + throw error(CL_INVALID_VALUE); + + // Create a temporary soft event that depends on ev, with + // pfn_event_notify as completion action. + ref_ptr<soft_event> sev = transfer( + new soft_event(ev->ctx, { ev }, true, + [=](event &) { + ev->wait(); + pfn_event_notify(ev, ev->status(), user_data); + })); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clRetainEvent(cl_event ev) { + if (!ev) + return CL_INVALID_EVENT; + + ev->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseEvent(cl_event ev) { + if (!ev) + return CL_INVALID_EVENT; + + if (ev->release()) + delete ev; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clEnqueueMarker(cl_command_queue q, cl_event *ev) try { + if (!q) + throw error(CL_INVALID_COMMAND_QUEUE); + + if (!ev) + throw error(CL_INVALID_VALUE); + + *ev = new hard_event(*q, CL_COMMAND_MARKER, {}); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueBarrier(cl_command_queue q) { + if (!q) + return CL_INVALID_COMMAND_QUEUE; + + // No need to do anything, q preserves data ordering strictly. + return CL_SUCCESS; +} + +PUBLIC cl_int +clEnqueueWaitForEvents(cl_command_queue q, cl_uint num_evs, + const cl_event *evs) try { + if (!q) + throw error(CL_INVALID_COMMAND_QUEUE); + + if (!num_evs || !evs) + throw error(CL_INVALID_VALUE); + + std::for_each(evs, evs + num_evs, [&](const cl_event ev) { + if (!ev) + throw error(CL_INVALID_EVENT); + + if (&ev->ctx != &q->ctx) + throw error(CL_INVALID_CONTEXT); + }); + + // Create a hard event that depends on the events in the wait list: + // subsequent commands in the same queue will be implicitly + // serialized with respect to it -- hard events always are. + ref_ptr<hard_event> hev = transfer( + new hard_event(*q, 0, { evs, evs + num_evs })); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clGetEventProfilingInfo(cl_event ev, cl_profiling_info param, + size_t size, void *buf, size_t *size_ret) { + return CL_PROFILING_INFO_NOT_AVAILABLE; +} + +PUBLIC cl_int +clFinish(cl_command_queue q) try { + if (!q) + throw error(CL_INVALID_COMMAND_QUEUE); + + // Create a temporary hard event -- it implicitly depends on all + // the previously queued hard events. + ref_ptr<hard_event> hev = transfer(new hard_event(*q, 0, { })); + + // And wait on it. + hev->wait(); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp new file mode 100644 index 00000000000..44eeb277127 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/kernel.cpp @@ -0,0 +1,318 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/kernel.hpp" +#include "core/event.hpp" + +using namespace clover; + +PUBLIC cl_kernel +clCreateKernel(cl_program prog, const char *name, + cl_int *errcode_ret) try { + if (!prog) + throw error(CL_INVALID_PROGRAM); + + if (!name) + throw error(CL_INVALID_VALUE); + + if (prog->binaries().empty()) + throw error(CL_INVALID_PROGRAM_EXECUTABLE); + + auto sym = prog->binaries().begin()->second.sym(name); + + ret_error(errcode_ret, CL_SUCCESS); + return new kernel(*prog, name, { sym.args.begin(), sym.args.end() }); + +} catch (module::noent_error &e) { + ret_error(errcode_ret, CL_INVALID_KERNEL_NAME); + return NULL; + +} catch(error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clCreateKernelsInProgram(cl_program prog, cl_uint count, + cl_kernel *kerns, cl_uint *count_ret) { + if (!prog) + throw error(CL_INVALID_PROGRAM); + + if (prog->binaries().empty()) + throw error(CL_INVALID_PROGRAM_EXECUTABLE); + + auto &syms = prog->binaries().begin()->second.syms; + + if (kerns && count < syms.size()) + throw error(CL_INVALID_VALUE); + + if (kerns) + std::transform(syms.begin(), syms.end(), kerns, + [=](const module::symbol &sym) { + return new kernel(*prog, compat::string(sym.name), + { sym.args.begin(), sym.args.end() }); + }); + + if (count_ret) + *count_ret = syms.size(); + + return CL_SUCCESS; +} + +PUBLIC cl_int +clRetainKernel(cl_kernel kern) { + if (!kern) + return CL_INVALID_KERNEL; + + kern->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseKernel(cl_kernel kern) { + if (!kern) + return CL_INVALID_KERNEL; + + if (kern->release()) + delete kern; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size, + const void *value) try { + if (!kern) + throw error(CL_INVALID_KERNEL); + + if (idx >= kern->args.size()) + throw error(CL_INVALID_ARG_INDEX); + + kern->args[idx]->set(size, value); + + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clGetKernelInfo(cl_kernel kern, cl_kernel_info param, + size_t size, void *buf, size_t *size_ret) { + if (!kern) + return CL_INVALID_KERNEL; + + switch (param) { + case CL_KERNEL_FUNCTION_NAME: + return string_property(buf, size, size_ret, kern->name()); + + case CL_KERNEL_NUM_ARGS: + return scalar_property<cl_uint>(buf, size, size_ret, + kern->args.size()); + + case CL_KERNEL_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, + kern->ref_count()); + + case CL_KERNEL_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, + &kern->prog.ctx); + + case CL_KERNEL_PROGRAM: + return scalar_property<cl_program>(buf, size, size_ret, + &kern->prog); + + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev, + cl_kernel_work_group_info param, + size_t size, void *buf, size_t *size_ret) { + if (!kern) + return CL_INVALID_KERNEL; + + if ((!dev && kern->prog.binaries().size() != 1) || + (dev && !kern->prog.binaries().count(dev))) + return CL_INVALID_DEVICE; + + switch (param) { + case CL_KERNEL_WORK_GROUP_SIZE: + return scalar_property<size_t>(buf, size, size_ret, + kern->max_block_size()); + + case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: + return vector_property<size_t>(buf, size, size_ret, + kern->block_size()); + + case CL_KERNEL_LOCAL_MEM_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, + kern->mem_local()); + + case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: + return scalar_property<size_t>(buf, size, size_ret, 1); + + case CL_KERNEL_PRIVATE_MEM_SIZE: + return scalar_property<cl_ulong>(buf, size, size_ret, + kern->mem_private()); + + default: + return CL_INVALID_VALUE; + } +} + +namespace { + /// + /// Common argument checking shared by kernel invocation commands. + /// + void + kernel_validate(cl_command_queue q, cl_kernel kern, + cl_uint dims, const size_t *grid_offset, + const size_t *grid_size, const size_t *block_size, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) { + if (!q) + throw error(CL_INVALID_COMMAND_QUEUE); + + if (!kern) + throw error(CL_INVALID_KERNEL); + + if (&kern->prog.ctx != &q->ctx || + any_of([&](const cl_event ev) { + return &ev->ctx != &q->ctx; + }, deps, deps + num_deps)) + throw error(CL_INVALID_CONTEXT); + + if (bool(num_deps) != bool(deps) || + any_of(is_zero<cl_event>(), deps, deps + num_deps)) + throw error(CL_INVALID_EVENT_WAIT_LIST); + + if (any_of([](std::unique_ptr<kernel::argument> &arg) { + return !arg->set(); + }, kern->args.begin(), kern->args.end())) + throw error(CL_INVALID_KERNEL_ARGS); + + if (!kern->prog.binaries().count(&q->dev)) + throw error(CL_INVALID_PROGRAM_EXECUTABLE); + + if (dims < 1 || dims > q->dev.max_block_size().size()) + throw error(CL_INVALID_WORK_DIMENSION); + + if (!grid_size || any_of(is_zero<size_t>(), grid_size, grid_size + dims)) + throw error(CL_INVALID_GLOBAL_WORK_SIZE); + + if (block_size && any_of([](size_t b, size_t max) { + return b == 0 || b > max; + }, block_size, block_size + dims, + q->dev.max_block_size().begin())) + throw error(CL_INVALID_WORK_ITEM_SIZE); + + if (block_size && any_of([](size_t b, size_t g) { + return g % b; + }, block_size, block_size + dims, grid_size)) + throw error(CL_INVALID_WORK_GROUP_SIZE); + } + + /// + /// Common event action shared by kernel invocation commands. + /// + std::function<void (event &)> + kernel_op(cl_command_queue q, cl_kernel kern, + const std::vector<size_t> &grid_offset, + const std::vector<size_t> &grid_size, + const std::vector<size_t> &block_size) { + const std::vector<size_t> reduced_grid_size = map( + std::divides<size_t>(), grid_size.begin(), grid_size.end(), + block_size.begin()); + + return [=](event &) { + kern->launch(*q, grid_offset, reduced_grid_size, block_size); + }; + } + + template<typename T, typename S> + std::vector<T> + opt_vector(const T *p, S n) { + if (p) + return { p, p + n }; + else + return { n }; + } +} + +PUBLIC cl_int +clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern, + cl_uint dims, const size_t *pgrid_offset, + const size_t *pgrid_size, const size_t *pblock_size, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + const std::vector<size_t> grid_offset = opt_vector(pgrid_offset, dims); + const std::vector<size_t> grid_size = opt_vector(pgrid_size, dims); + const std::vector<size_t> block_size = opt_vector(pblock_size, dims); + + kernel_validate(q, kern, dims, pgrid_offset, pgrid_size, pblock_size, + num_deps, deps, ev); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_NDRANGE_KERNEL, { deps, deps + num_deps }, + kernel_op(q, kern, grid_offset, grid_size, block_size)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueTask(cl_command_queue q, cl_kernel kern, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + const std::vector<size_t> grid_offset = { 0 }; + const std::vector<size_t> grid_size = { 1 }; + const std::vector<size_t> block_size = { 1 }; + + kernel_validate(q, kern, 1, grid_offset.data(), grid_size.data(), + block_size.data(), num_deps, deps, ev); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_TASK, { deps, deps + num_deps }, + kernel_op(q, kern, grid_offset, grid_size, block_size)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch(error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueNativeKernel(cl_command_queue q, void (*func)(void *), + void *args, size_t args_size, + cl_uint obj_count, const cl_mem *obj_list, + const void **obj_args, cl_uint num_deps, + const cl_event *deps, cl_event *ev) { + return CL_INVALID_OPERATION; +} diff --git a/src/gallium/state_trackers/clover/api/memory.cpp b/src/gallium/state_trackers/clover/api/memory.cpp new file mode 100644 index 00000000000..1b1ae73796f --- /dev/null +++ b/src/gallium/state_trackers/clover/api/memory.cpp @@ -0,0 +1,305 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/memory.hpp" +#include "core/format.hpp" + +using namespace clover; + +PUBLIC cl_mem +clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size, + void *host_ptr, cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_COPY_HOST_PTR))) + throw error(CL_INVALID_HOST_PTR); + + if (!size) + throw error(CL_INVALID_BUFFER_SIZE); + + if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | + CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | + CL_MEM_COPY_HOST_PTR)) + throw error(CL_INVALID_VALUE); + + ret_error(errcode_ret, CL_SUCCESS); + return new root_buffer(*ctx, flags, size, host_ptr); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_mem +clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op, + const void *op_info, cl_int *errcode_ret) try { + root_buffer *parent = dynamic_cast<root_buffer *>(obj); + + if (!parent) + throw error(CL_INVALID_MEM_OBJECT); + + if ((flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_ALLOC_HOST_PTR | + CL_MEM_COPY_HOST_PTR)) || + (~flags & parent->flags() & (CL_MEM_READ_ONLY | + CL_MEM_WRITE_ONLY))) + throw error(CL_INVALID_VALUE); + + if (op == CL_BUFFER_CREATE_TYPE_REGION) { + const cl_buffer_region *reg = (const cl_buffer_region *)op_info; + + if (!reg || + reg->origin > parent->size() || + reg->origin + reg->size > parent->size()) + throw error(CL_INVALID_VALUE); + + if (!reg->size) + throw error(CL_INVALID_BUFFER_SIZE); + + ret_error(errcode_ret, CL_SUCCESS); + return new sub_buffer(*parent, flags, reg->origin, reg->size); + + } else { + throw error(CL_INVALID_VALUE); + } + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_mem +clCreateImage2D(cl_context ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t row_pitch, + void *host_ptr, cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | + CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | + CL_MEM_COPY_HOST_PTR)) + throw error(CL_INVALID_VALUE); + + if (!format) + throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR); + + if (width < 1 || height < 1) + throw error(CL_INVALID_IMAGE_SIZE); + + if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_COPY_HOST_PTR))) + throw error(CL_INVALID_HOST_PTR); + + if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE2D).count(*format)) + throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED); + + ret_error(errcode_ret, CL_SUCCESS); + return new image2d(*ctx, flags, format, width, height, + row_pitch, host_ptr); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_mem +clCreateImage3D(cl_context ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, + void *host_ptr, cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | + CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | + CL_MEM_COPY_HOST_PTR)) + throw error(CL_INVALID_VALUE); + + if (!format) + throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR); + + if (width < 1 || height < 1 || depth < 2) + throw error(CL_INVALID_IMAGE_SIZE); + + if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR | + CL_MEM_COPY_HOST_PTR))) + throw error(CL_INVALID_HOST_PTR); + + if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE3D).count(*format)) + throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED); + + ret_error(errcode_ret, CL_SUCCESS); + return new image3d(*ctx, flags, format, width, height, depth, + row_pitch, slice_pitch, host_ptr); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clGetSupportedImageFormats(cl_context ctx, cl_mem_flags flags, + cl_mem_object_type type, cl_uint count, + cl_image_format *buf, cl_uint *count_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY | + CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | + CL_MEM_COPY_HOST_PTR)) + throw error(CL_INVALID_VALUE); + + if (!count && buf) + throw error(CL_INVALID_VALUE); + + auto formats = supported_formats(ctx, type); + + if (buf) + std::copy_n(formats.begin(), std::min((cl_uint)formats.size(), count), + buf); + if (count_ret) + *count_ret = formats.size(); + + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clGetMemObjectInfo(cl_mem obj, cl_mem_info param, + size_t size, void *buf, size_t *size_ret) { + if (!obj) + return CL_INVALID_MEM_OBJECT; + + switch (param) { + case CL_MEM_TYPE: + return scalar_property<cl_mem_object_type>(buf, size, size_ret, + obj->type()); + + case CL_MEM_FLAGS: + return scalar_property<cl_mem_flags>(buf, size, size_ret, obj->flags()); + + case CL_MEM_SIZE: + return scalar_property<size_t>(buf, size, size_ret, obj->size()); + + case CL_MEM_HOST_PTR: + return scalar_property<void *>(buf, size, size_ret, obj->host_ptr()); + + case CL_MEM_MAP_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, 0); + + case CL_MEM_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, obj->ref_count()); + + case CL_MEM_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, &obj->ctx); + + case CL_MEM_ASSOCIATED_MEMOBJECT: { + sub_buffer *sub = dynamic_cast<sub_buffer *>(obj); + return scalar_property<cl_mem>(buf, size, size_ret, + (sub ? &sub->parent : NULL)); + } + case CL_MEM_OFFSET: { + sub_buffer *sub = dynamic_cast<sub_buffer *>(obj); + return scalar_property<size_t>(buf, size, size_ret, + (sub ? sub->offset() : 0)); + } + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clGetImageInfo(cl_mem obj, cl_image_info param, + size_t size, void *buf, size_t *size_ret) { + image *img = dynamic_cast<image *>(obj); + if (!img) + return CL_INVALID_MEM_OBJECT; + + switch (param) { + case CL_IMAGE_FORMAT: + return scalar_property<cl_image_format>(buf, size, size_ret, + img->format()); + + case CL_IMAGE_ELEMENT_SIZE: + return scalar_property<size_t>(buf, size, size_ret, 0); + + case CL_IMAGE_ROW_PITCH: + return scalar_property<size_t>(buf, size, size_ret, img->row_pitch()); + + case CL_IMAGE_SLICE_PITCH: + return scalar_property<size_t>(buf, size, size_ret, img->slice_pitch()); + + case CL_IMAGE_WIDTH: + return scalar_property<size_t>(buf, size, size_ret, img->width()); + + case CL_IMAGE_HEIGHT: + return scalar_property<size_t>(buf, size, size_ret, img->height()); + + case CL_IMAGE_DEPTH: + return scalar_property<size_t>(buf, size, size_ret, img->depth()); + + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clRetainMemObject(cl_mem obj) { + if (!obj) + return CL_INVALID_MEM_OBJECT; + + obj->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseMemObject(cl_mem obj) { + if (!obj) + return CL_INVALID_MEM_OBJECT; + + if (obj->release()) + delete obj; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clSetMemObjectDestructorCallback(cl_mem obj, + void (CL_CALLBACK *pfn_notify)(cl_mem, void *), + void *user_data) { + if (!obj) + return CL_INVALID_MEM_OBJECT; + + if (!pfn_notify) + return CL_INVALID_VALUE; + + obj->destroy_notify([=]{ pfn_notify(obj, user_data); }); + + return CL_SUCCESS; +} diff --git a/src/gallium/state_trackers/clover/api/platform.cpp b/src/gallium/state_trackers/clover/api/platform.cpp new file mode 100644 index 00000000000..e5e80b85256 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/platform.cpp @@ -0,0 +1,68 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" + +using namespace clover; + +PUBLIC cl_int +clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, + cl_uint *num_platforms) { + if ((!num_entries && platforms) || + (!num_platforms && !platforms)) + return CL_INVALID_VALUE; + + if (num_platforms) + *num_platforms = 1; + if (platforms) + *platforms = NULL; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name, + size_t size, void *buf, size_t *size_ret) { + if (platform != NULL) + return CL_INVALID_PLATFORM; + + switch (param_name) { + case CL_PLATFORM_PROFILE: + return string_property(buf, size, size_ret, "FULL_PROFILE"); + + case CL_PLATFORM_VERSION: + return string_property(buf, size, size_ret, + "OpenCL 1.1 MESA " MESA_VERSION); + + case CL_PLATFORM_NAME: + return string_property(buf, size, size_ret, "Default"); + + case CL_PLATFORM_VENDOR: + return string_property(buf, size, size_ret, "Mesa"); + + case CL_PLATFORM_EXTENSIONS: + return string_property(buf, size, size_ret, ""); + + default: + return CL_INVALID_VALUE; + } +} diff --git a/src/gallium/state_trackers/clover/api/program.cpp b/src/gallium/state_trackers/clover/api/program.cpp new file mode 100644 index 00000000000..e874c51ad7d --- /dev/null +++ b/src/gallium/state_trackers/clover/api/program.cpp @@ -0,0 +1,241 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/program.hpp" + +using namespace clover; + +PUBLIC cl_program +clCreateProgramWithSource(cl_context ctx, cl_uint count, + const char **strings, const size_t *lengths, + cl_int *errcode_ret) try { + std::string source; + + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (!count || !strings || + any_of(is_zero<const char *>(), strings, strings + count)) + throw error(CL_INVALID_VALUE); + + // Concatenate all the provided fragments together + for (unsigned i = 0; i < count; ++i) + source += (lengths && lengths[i] ? + std::string(strings[i], strings[i] + lengths[i]) : + std::string(strings[i])); + + // ...and create a program object for them. + ret_error(errcode_ret, CL_SUCCESS); + return new program(*ctx, source); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_program +clCreateProgramWithBinary(cl_context ctx, cl_uint count, + const cl_device_id *devs, const size_t *lengths, + const unsigned char **binaries, cl_int *status_ret, + cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (!count || !devs || !lengths || !binaries) + throw error(CL_INVALID_VALUE); + + if (any_of([&](const cl_device_id dev) { + return !ctx->has_device(dev); + }, devs, devs + count)) + throw error(CL_INVALID_DEVICE); + + // Deserialize the provided binaries, + auto modules = map( + [](const unsigned char *p, size_t l) -> std::pair<cl_int, module> { + if (!p || !l) + return { CL_INVALID_VALUE, {} }; + + try { + compat::istream::buffer_t bin(p, l); + compat::istream s(bin); + + return { CL_SUCCESS, module::deserialize(s) }; + + } catch (compat::istream::error &e) { + return { CL_INVALID_BINARY, {} }; + } + }, + binaries, binaries + count, lengths); + + // update the status array, + if (status_ret) + std::transform(modules.begin(), modules.end(), status_ret, + keys<cl_int, module>); + + if (any_of(key_equals<cl_int, module>(CL_INVALID_VALUE), + modules.begin(), modules.end())) + throw error(CL_INVALID_VALUE); + + if (any_of(key_equals<cl_int, module>(CL_INVALID_BINARY), + modules.begin(), modules.end())) + throw error(CL_INVALID_BINARY); + + // initialize a program object with them. + ret_error(errcode_ret, CL_SUCCESS); + return new program(*ctx, { devs, devs + count }, + map(values<cl_int, module>, + modules.begin(), modules.end())); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clRetainProgram(cl_program prog) { + if (!prog) + return CL_INVALID_PROGRAM; + + prog->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseProgram(cl_program prog) { + if (!prog) + return CL_INVALID_PROGRAM; + + if (prog->release()) + delete prog; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clBuildProgram(cl_program prog, cl_uint count, const cl_device_id *devs, + const char *opts, void (*pfn_notify)(cl_program, void *), + void *user_data) try { + if (!prog) + throw error(CL_INVALID_PROGRAM); + + if (bool(count) != bool(devs) || + (!pfn_notify && user_data)) + throw error(CL_INVALID_VALUE); + + if (any_of([&](const cl_device_id dev) { + return !prog->ctx.has_device(dev); + }, devs, devs + count)) + throw error(CL_INVALID_DEVICE); + + prog->build({ devs, devs + count }); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clUnloadCompiler() { + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetProgramInfo(cl_program prog, cl_program_info param, + size_t size, void *buf, size_t *size_ret) { + if (!prog) + return CL_INVALID_PROGRAM; + + switch (param) { + case CL_PROGRAM_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, + prog->ref_count()); + + case CL_PROGRAM_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, + &prog->ctx); + + case CL_PROGRAM_NUM_DEVICES: + return scalar_property<cl_uint>(buf, size, size_ret, + prog->binaries().size()); + + case CL_PROGRAM_DEVICES: + return vector_property<cl_device_id>( + buf, size, size_ret, + map(keys<device *, module>, + prog->binaries().begin(), prog->binaries().end())); + + case CL_PROGRAM_SOURCE: + return string_property(buf, size, size_ret, prog->source()); + + case CL_PROGRAM_BINARY_SIZES: + return vector_property<size_t>( + buf, size, size_ret, + map([](const std::pair<device *, module> &ent) { + compat::ostream::buffer_t bin; + compat::ostream s(bin); + ent.second.serialize(s); + return bin.size(); + }, + prog->binaries().begin(), prog->binaries().end())); + + case CL_PROGRAM_BINARIES: + return matrix_property<unsigned char>( + buf, size, size_ret, + map([](const std::pair<device *, module> &ent) { + compat::ostream::buffer_t bin; + compat::ostream s(bin); + ent.second.serialize(s); + return bin; + }, + prog->binaries().begin(), prog->binaries().end())); + + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clGetProgramBuildInfo(cl_program prog, cl_device_id dev, + cl_program_build_info param, + size_t size, void *buf, size_t *size_ret) { + if (!prog) + return CL_INVALID_PROGRAM; + + if (!prog->ctx.has_device(dev)) + return CL_INVALID_DEVICE; + + switch (param) { + case CL_PROGRAM_BUILD_STATUS: + return scalar_property<cl_build_status>(buf, size, size_ret, + prog->build_status(dev)); + + case CL_PROGRAM_BUILD_OPTIONS: + return string_property(buf, size, size_ret, prog->build_opts(dev)); + + case CL_PROGRAM_BUILD_LOG: + return string_property(buf, size, size_ret, prog->build_log(dev)); + + default: + return CL_INVALID_VALUE; + } +} diff --git a/src/gallium/state_trackers/clover/api/queue.cpp b/src/gallium/state_trackers/clover/api/queue.cpp new file mode 100644 index 00000000000..a7905bc4396 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/queue.cpp @@ -0,0 +1,102 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/queue.hpp" + +using namespace clover; + +PUBLIC cl_command_queue +clCreateCommandQueue(cl_context ctx, cl_device_id dev, + cl_command_queue_properties props, + cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + if (!ctx->has_device(dev)) + throw error(CL_INVALID_DEVICE); + + if (props & ~(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | + CL_QUEUE_PROFILING_ENABLE)) + throw error(CL_INVALID_VALUE); + + ret_error(errcode_ret, CL_SUCCESS); + return new command_queue(*ctx, *dev, props); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clRetainCommandQueue(cl_command_queue q) { + if (!q) + return CL_INVALID_COMMAND_QUEUE; + + q->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseCommandQueue(cl_command_queue q) { + if (!q) + return CL_INVALID_COMMAND_QUEUE; + + if (q->release()) + delete q; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param, + size_t size, void *buf, size_t *size_ret) { + if (!q) + return CL_INVALID_COMMAND_QUEUE; + + switch (param) { + case CL_QUEUE_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, &q->ctx); + + case CL_QUEUE_DEVICE: + return scalar_property<cl_device_id>(buf, size, size_ret, &q->dev); + + case CL_QUEUE_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, q->ref_count()); + + case CL_QUEUE_PROPERTIES: + return scalar_property<cl_command_queue_properties>(buf, size, size_ret, + q->props()); + + default: + return CL_INVALID_VALUE; + } +} + +PUBLIC cl_int +clFlush(cl_command_queue q) { + if (!q) + return CL_INVALID_COMMAND_QUEUE; + + q->flush(); + return CL_SUCCESS; +} diff --git a/src/gallium/state_trackers/clover/api/sampler.cpp b/src/gallium/state_trackers/clover/api/sampler.cpp new file mode 100644 index 00000000000..32ce22ef90f --- /dev/null +++ b/src/gallium/state_trackers/clover/api/sampler.cpp @@ -0,0 +1,90 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "api/util.hpp" +#include "core/sampler.hpp" + +using namespace clover; + +PUBLIC cl_sampler +clCreateSampler(cl_context ctx, cl_bool norm_mode, + cl_addressing_mode addr_mode, cl_filter_mode filter_mode, + cl_int *errcode_ret) try { + if (!ctx) + throw error(CL_INVALID_CONTEXT); + + ret_error(errcode_ret, CL_SUCCESS); + return new sampler(*ctx, norm_mode, addr_mode, filter_mode); + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clRetainSampler(cl_sampler s) { + if (!s) + throw error(CL_INVALID_SAMPLER); + + s->retain(); + return CL_SUCCESS; +} + +PUBLIC cl_int +clReleaseSampler(cl_sampler s) { + if (!s) + throw error(CL_INVALID_SAMPLER); + + if (s->release()) + delete s; + + return CL_SUCCESS; +} + +PUBLIC cl_int +clGetSamplerInfo(cl_sampler s, cl_sampler_info param, + size_t size, void *buf, size_t *size_ret) { + if (!s) + throw error(CL_INVALID_SAMPLER); + + switch (param) { + case CL_SAMPLER_REFERENCE_COUNT: + return scalar_property<cl_uint>(buf, size, size_ret, s->ref_count()); + + case CL_SAMPLER_CONTEXT: + return scalar_property<cl_context>(buf, size, size_ret, &s->ctx); + + case CL_SAMPLER_NORMALIZED_COORDS: + return scalar_property<cl_bool>(buf, size, size_ret, s->norm_mode()); + + case CL_SAMPLER_ADDRESSING_MODE: + return scalar_property<cl_addressing_mode>(buf, size, size_ret, + s->addr_mode()); + + case CL_SAMPLER_FILTER_MODE: + return scalar_property<cl_filter_mode>(buf, size, size_ret, + s->filter_mode()); + + default: + return CL_INVALID_VALUE; + } +} diff --git a/src/gallium/state_trackers/clover/api/transfer.cpp b/src/gallium/state_trackers/clover/api/transfer.cpp new file mode 100644 index 00000000000..c67b75e8034 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/transfer.cpp @@ -0,0 +1,506 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <cstring> + +#include "api/util.hpp" +#include "core/event.hpp" +#include "core/resource.hpp" + +using namespace clover; + +namespace { + typedef resource::point point; + + /// + /// Common argument checking shared by memory transfer commands. + /// + void + validate_base(cl_command_queue q, cl_uint num_deps, const cl_event *deps) { + if (!q) + throw error(CL_INVALID_COMMAND_QUEUE); + + if (bool(num_deps) != bool(deps) || + any_of(is_zero<cl_event>(), deps, deps + num_deps)) + throw error(CL_INVALID_EVENT_WAIT_LIST); + + if (any_of([&](const cl_event ev) { + return &ev->ctx != &q->ctx; + }, deps, deps + num_deps)) + throw error(CL_INVALID_CONTEXT); + } + + /// + /// Memory object-specific argument checking shared by most memory + /// transfer commands. + /// + void + validate_obj(cl_command_queue q, cl_mem obj) { + if (!obj) + throw error(CL_INVALID_MEM_OBJECT); + + if (&obj->ctx != &q->ctx) + throw error(CL_INVALID_CONTEXT); + } + + /// + /// Class that encapsulates the task of mapping an object of type + /// \a T. The return value of get() should be implicitly + /// convertible to \a void *. + /// + template<typename T> struct __map; + + template<> struct __map<void *> { + static void * + get(cl_command_queue q, void *obj, cl_map_flags flags, + size_t offset, size_t size) { + return (char *)obj + offset; + } + }; + + template<> struct __map<const void *> { + static const void * + get(cl_command_queue q, const void *obj, cl_map_flags flags, + size_t offset, size_t size) { + return (const char *)obj + offset; + } + }; + + template<> struct __map<memory_obj *> { + static mapping + get(cl_command_queue q, memory_obj *obj, cl_map_flags flags, + size_t offset, size_t size) { + return { *q, obj->resource(q), flags, true, { offset }, { size }}; + } + }; + + /// + /// Software copy from \a src_obj to \a dst_obj. They can be + /// either pointers or memory objects. + /// + template<typename T, typename S> + std::function<void (event &)> + soft_copy_op(cl_command_queue q, + T dst_obj, const point &dst_orig, const point &dst_pitch, + S src_obj, const point &src_orig, const point &src_pitch, + const point ®ion) { + return [=](event &) { + auto dst = __map<T>::get(q, dst_obj, CL_MAP_WRITE, + dst_pitch(dst_orig), dst_pitch(region)); + auto src = __map<S>::get(q, src_obj, CL_MAP_READ, + src_pitch(src_orig), src_pitch(region)); + point p; + + for (p[2] = 0; p[2] < region[2]; ++p[2]) { + for (p[1] = 0; p[1] < region[1]; ++p[1]) { + std::memcpy(static_cast<char *>(dst) + dst_pitch(p), + static_cast<const char *>(src) + src_pitch(p), + src_pitch[0] * region[0]); + } + } + }; + } + + /// + /// Hardware copy from \a src_obj to \a dst_obj. + /// + template<typename T, typename S> + std::function<void (event &)> + hard_copy_op(cl_command_queue q, T dst_obj, const point &dst_orig, + S src_obj, const point &src_orig, const point ®ion) { + return [=](event &) { + dst_obj->resource(q).copy(*q, dst_orig, region, + src_obj->resource(q), src_orig); + }; + } +} + +PUBLIC cl_int +clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, + size_t offset, size_t size, void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + if (!ptr || offset > obj->size() || offset + size > obj->size()) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_READ_BUFFER, { deps, deps + num_deps }, + soft_copy_op(q, + ptr, { 0 }, { 1 }, + obj, { offset }, { 1 }, + { size, 1, 1 })); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, + size_t offset, size_t size, const void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + if (!ptr || offset > obj->size() || offset + size > obj->size()) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_WRITE_BUFFER, { deps, deps + num_deps }, + soft_copy_op(q, + obj, { offset }, { 1 }, + ptr, { 0 }, { 1 }, + { size, 1, 1 })); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, + const size_t *obj_origin, const size_t *host_origin, + const size_t *region, + size_t obj_row_pitch, size_t obj_slice_pitch, + size_t host_row_pitch, size_t host_slice_pitch, + void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + if (!ptr) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_READ_BUFFER_RECT, { deps, deps + num_deps }, + soft_copy_op(q, + ptr, host_origin, + { 1, host_row_pitch, host_slice_pitch }, + obj, obj_origin, + { 1, obj_row_pitch, obj_slice_pitch }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking, + const size_t *obj_origin, const size_t *host_origin, + const size_t *region, + size_t obj_row_pitch, size_t obj_slice_pitch, + size_t host_row_pitch, size_t host_slice_pitch, + const void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + if (!ptr) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_WRITE_BUFFER_RECT, { deps, deps + num_deps }, + soft_copy_op(q, + obj, obj_origin, + { 1, obj_row_pitch, obj_slice_pitch }, + ptr, host_origin, + { 1, host_row_pitch, host_slice_pitch }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, + size_t src_offset, size_t dst_offset, size_t size, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, src_obj); + validate_obj(q, dst_obj); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_COPY_BUFFER, { deps, deps + num_deps }, + hard_copy_op(q, dst_obj, { dst_offset }, + src_obj, { src_offset }, + { size, 1, 1 })); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, + const size_t *src_origin, const size_t *dst_origin, + const size_t *region, + size_t src_row_pitch, size_t src_slice_pitch, + size_t dst_row_pitch, size_t dst_slice_pitch, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, src_obj); + validate_obj(q, dst_obj); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_COPY_BUFFER_RECT, { deps, deps + num_deps }, + soft_copy_op(q, + dst_obj, dst_origin, + { 1, dst_row_pitch, dst_slice_pitch }, + src_obj, src_origin, + { 1, src_row_pitch, src_slice_pitch }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking, + const size_t *origin, const size_t *region, + size_t row_pitch, size_t slice_pitch, void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + image *img = dynamic_cast<image *>(obj); + + validate_base(q, num_deps, deps); + validate_obj(q, img); + + if (!ptr) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_READ_IMAGE, { deps, deps + num_deps }, + soft_copy_op(q, + ptr, {}, + { 1, row_pitch, slice_pitch }, + obj, origin, + { 1, img->row_pitch(), img->slice_pitch() }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking, + const size_t *origin, const size_t *region, + size_t row_pitch, size_t slice_pitch, const void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + image *img = dynamic_cast<image *>(obj); + + validate_base(q, num_deps, deps); + validate_obj(q, img); + + if (!ptr) + throw error(CL_INVALID_VALUE); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_WRITE_IMAGE, { deps, deps + num_deps }, + soft_copy_op(q, + obj, origin, + { 1, img->row_pitch(), img->slice_pitch() }, + ptr, {}, + { 1, row_pitch, slice_pitch }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, + const size_t *src_origin, const size_t *dst_origin, + const size_t *region, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + image *src_img = dynamic_cast<image *>(src_obj); + image *dst_img = dynamic_cast<image *>(dst_obj); + + validate_base(q, num_deps, deps); + validate_obj(q, src_img); + validate_obj(q, dst_img); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_COPY_IMAGE, { deps, deps + num_deps }, + hard_copy_op(q, dst_obj, dst_origin, src_obj, src_origin, region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, + const size_t *src_origin, const size_t *region, + size_t dst_offset, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + image *src_img = dynamic_cast<image *>(src_obj); + + validate_base(q, num_deps, deps); + validate_obj(q, src_img); + validate_obj(q, dst_obj); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, { deps, deps + num_deps }, + soft_copy_op(q, + dst_obj, { dst_offset }, + { 0, 0, 0 }, + src_obj, src_origin, + { 1, src_img->row_pitch(), src_img->slice_pitch() }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC cl_int +clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj, + size_t src_offset, + const size_t *dst_origin, const size_t *region, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + image *dst_img = dynamic_cast<image *>(src_obj); + + validate_base(q, num_deps, deps); + validate_obj(q, src_obj); + validate_obj(q, dst_img); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, { deps, deps + num_deps }, + soft_copy_op(q, + dst_obj, dst_origin, + { 1, dst_img->row_pitch(), dst_img->slice_pitch() }, + src_obj, { src_offset }, + { 0, 0, 0 }, + region)); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} + +PUBLIC void * +clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking, + cl_map_flags flags, size_t offset, size_t size, + cl_uint num_deps, const cl_event *deps, + cl_event *ev, cl_int *errcode_ret) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + if (offset > obj->size() || offset + size > obj->size()) + throw error(CL_INVALID_VALUE); + + void *map = obj->resource(q).add_map( + *q, flags, blocking, { offset }, { size }); + + ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER, + { deps, deps + num_deps })); + ret_error(errcode_ret, CL_SUCCESS); + return map; + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC void * +clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking, + cl_map_flags flags, + const size_t *origin, const size_t *region, + size_t *row_pitch, size_t *slice_pitch, + cl_uint num_deps, const cl_event *deps, + cl_event *ev, cl_int *errcode_ret) try { + image *img = dynamic_cast<image *>(obj); + + validate_base(q, num_deps, deps); + validate_obj(q, img); + + void *map = obj->resource(q).add_map( + *q, flags, blocking, origin, region); + + ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE, + { deps, deps + num_deps })); + ret_error(errcode_ret, CL_SUCCESS); + return map; + +} catch (error &e) { + ret_error(errcode_ret, e); + return NULL; +} + +PUBLIC cl_int +clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr, + cl_uint num_deps, const cl_event *deps, + cl_event *ev) try { + validate_base(q, num_deps, deps); + validate_obj(q, obj); + + hard_event *hev = new hard_event( + *q, CL_COMMAND_UNMAP_MEM_OBJECT, { deps, deps + num_deps }, + [=](event &) { + obj->resource(q).del_map(ptr); + }); + + ret_object(ev, hev); + return CL_SUCCESS; + +} catch (error &e) { + return e.get(); +} diff --git a/src/gallium/state_trackers/clover/api/util.hpp b/src/gallium/state_trackers/clover/api/util.hpp new file mode 100644 index 00000000000..2f9ec1f6a10 --- /dev/null +++ b/src/gallium/state_trackers/clover/api/util.hpp @@ -0,0 +1,166 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CL_UTIL_HPP__ +#define __CL_UTIL_HPP__ + +#include <cstdint> +#include <cstring> +#include <algorithm> +#include <map> + +#include "core/base.hpp" +#include "pipe/p_compiler.h" + +namespace clover { + /// + /// Return a matrix (a container of containers) in \a buf with + /// argument and bounds checking. Intended to be used by + /// implementations of \a clGetXXXInfo(). + /// + template<typename T, typename V> + cl_int + matrix_property(void *buf, size_t size, size_t *size_ret, const V& v) { + if (buf && size < sizeof(T *) * v.size()) + return CL_INVALID_VALUE; + + if (size_ret) + *size_ret = sizeof(T *) * v.size(); + + if (buf) + for_each([](typename V::value_type src, T *dst) { + if (dst) + std::copy(src.begin(), src.end(), dst); + }, + v.begin(), v.end(), (T **)buf); + + return CL_SUCCESS; + } + + /// + /// Return a vector in \a buf with argument and bounds checking. + /// Intended to be used by implementations of \a clGetXXXInfo(). + /// + template<typename T, typename V> + cl_int + vector_property(void *buf, size_t size, size_t *size_ret, const V& v) { + if (buf && size < sizeof(T) * v.size()) + return CL_INVALID_VALUE; + + if (size_ret) + *size_ret = sizeof(T) * v.size(); + if (buf) + std::copy(v.begin(), v.end(), (T *)buf); + + return CL_SUCCESS; + } + + /// + /// Return a scalar in \a buf with argument and bounds checking. + /// Intended to be used by implementations of \a clGetXXXInfo(). + /// + template<typename T> + cl_int + scalar_property(void *buf, size_t size, size_t *size_ret, T v) { + return vector_property<T>(buf, size, size_ret, std::vector<T>(1, v)); + } + + /// + /// Return a string in \a buf with argument and bounds checking. + /// Intended to be used by implementations of \a clGetXXXInfo(). + /// + inline cl_int + string_property(void *buf, size_t size, size_t *size_ret, + const std::string &v) { + if (buf && size < v.size() + 1) + return CL_INVALID_VALUE; + + if (size_ret) + *size_ret = v.size() + 1; + if (buf) + std::strcpy((char *)buf, v.c_str()); + + return CL_SUCCESS; + } + + /// + /// Convert a NULL-terminated property list into an std::map. + /// + template<typename T> + std::map<T, T> + property_map(const T *props) { + std::map<T, T> m; + + while (props && *props) { + T key = *props++; + T value = *props++; + + if (m.count(key)) + throw clover::error(CL_INVALID_PROPERTY); + + m.insert({ key, value }); + } + + return m; + } + + /// + /// Convert an std::map into a NULL-terminated property list. + /// + template<typename T> + std::vector<T> + property_vector(const std::map<T, T> &m) { + std::vector<T> v; + + for (auto &p : m) { + v.push_back(p.first); + v.push_back(p.second); + } + + v.push_back(0); + return v; + } + + /// + /// Return an error code in \a p if non-zero. + /// + inline void + ret_error(cl_int *p, const clover::error &e) { + if (p) + *p = e.get(); + } + + /// + /// Return a reference-counted object in \a p if non-zero. + /// Otherwise release object ownership. + /// + template<typename T, typename S> + void + ret_object(T p, S v) { + if (p) + *p = v; + else + v->release(); + } +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/base.hpp b/src/gallium/state_trackers/clover/core/base.hpp new file mode 100644 index 00000000000..19053f39235 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/base.hpp @@ -0,0 +1,285 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_BASE_HPP__ +#define __CORE_BASE_HPP__ + +#include <stdexcept> +#include <atomic> +#include <cassert> +#include <tuple> +#include <vector> +#include <functional> + +#include "CL/cl.h" + +/// +/// Main namespace of the CL state tracker. +/// +namespace clover { + /// + /// Class that represents an error that can be converted to an + /// OpenCL status code. + /// + class error : public std::runtime_error { + public: + error(cl_int code, std::string what = "") : + std::runtime_error(what), code(code) { + } + + cl_int get() const { + return code; + } + + protected: + cl_int code; + }; + + /// + /// Base class for objects that support reference counting. + /// + class ref_counter { + public: + ref_counter() : __ref_count(1) {} + + unsigned ref_count() { + return __ref_count; + } + + void retain() { + __ref_count++; + } + + bool release() { + return (--__ref_count) == 0; + } + + private: + std::atomic<unsigned> __ref_count; + }; + + /// + /// Intrusive smart pointer for objects that implement the + /// clover::ref_counter interface. + /// + template<typename T> + class ref_ptr { + public: + ref_ptr(T *q = NULL) : p(NULL) { + reset(q); + } + + template<typename S> + ref_ptr(const ref_ptr<S> &ref) : p(NULL) { + reset(ref.p); + } + + ~ref_ptr() { + reset(NULL); + } + + void reset(T *q = NULL) { + if (q) + q->retain(); + if (p && p->release()) + delete p; + p = q; + } + + ref_ptr &operator=(const ref_ptr &ref) { + reset(ref.p); + return *this; + } + + T *operator*() const { + return p; + } + + T *operator->() const { + return p; + } + + operator bool() const { + return p; + } + + private: + T *p; + }; + + /// + /// Transfer the caller's ownership of a reference-counted object + /// to a clover::ref_ptr smart pointer. + /// + template<typename T> + inline ref_ptr<T> + transfer(T *p) { + ref_ptr<T> ref { p }; + p->release(); + return ref; + } + + template<typename T, typename S, int N> + struct __iter_helper { + template<typename F, typename Its, typename... Args> + static T + step(F op, S state, Its its, Args... args) { + return __iter_helper<T, S, N - 1>::step( + op, state, its, *(std::get<N>(its)++), args...); + } + }; + + template<typename T, typename S> + struct __iter_helper<T, S, 0> { + template<typename F, typename Its, typename... Args> + static T + step(F op, S state, Its its, Args... args) { + return op(state, *(std::get<0>(its)++), args...); + } + }; + + struct __empty {}; + + template<typename T> + struct __iter_helper<T, __empty, 0> { + template<typename F, typename Its, typename... Args> + static T + step(F op, __empty state, Its its, Args... args) { + return op(*(std::get<0>(its)++), args...); + } + }; + + template<typename F, typename... Its> + struct __result_helper { + typedef typename std::remove_const< + typename std::result_of< + F (typename std::iterator_traits<Its>::value_type...) + >::type + >::type type; + }; + + /// + /// Iterate \a op on the result of zipping all the specified + /// iterators together. + /// + /// Similar to std::for_each, but it accepts functions of an + /// arbitrary number of arguments. + /// + template<typename F, typename It0, typename... Its> + F + for_each(F op, It0 it0, It0 end0, Its... its) { + while (it0 != end0) + __iter_helper<void, __empty, sizeof...(Its)>::step( + op, {}, std::tie(it0, its...)); + + return op; + } + + /// + /// Iterate \a op on the result of zipping all the specified + /// iterators together, storing return values in a new container. + /// + /// Similar to std::transform, but it accepts functions of an + /// arbitrary number of arguments and it doesn't have to be + /// provided with an output iterator. + /// + template<typename F, typename It0, typename... Its, + typename C = std::vector< + typename __result_helper<F, It0, Its...>::type>> + C + map(F op, It0 it0, It0 end0, Its... its) { + C c; + + while (it0 != end0) + c.push_back( + __iter_helper<typename C::value_type, __empty, sizeof...(Its)> + ::step(op, {}, std::tie(it0, its...))); + + return c; + } + + /// + /// Reduce the result of zipping all the specified iterators + /// together, using iterative application of \a op from left to + /// right. + /// + /// Similar to std::accumulate, but it accepts functions of an + /// arbitrary number of arguments. + /// + template<typename F, typename T, typename It0, typename... Its> + T + fold(F op, T a, It0 it0, It0 end0, Its... its) { + while (it0 != end0) + a = __iter_helper<T, T, sizeof...(Its)>::step( + op, a, std::tie(it0, its...)); + + return a; + } + + /// + /// Iterate \a op on the result of zipping the specified iterators + /// together, checking if any of the evaluations returns \a true. + /// + /// Similar to std::any_of, but it accepts functions of an + /// arbitrary number of arguments. + /// + template<typename F, typename It0, typename... Its> + bool + any_of(F op, It0 it0, It0 end0, Its... its) { + while (it0 != end0) + if (__iter_helper<bool, __empty, sizeof...(Its)>::step( + op, {}, std::tie(it0, its...))) + return true; + + return false; + } + + template<typename T, typename S> + T + keys(const std::pair<T, S> &ent) { + return ent.first; + } + + template<typename T, typename S> + std::function<bool (const std::pair<T, S> &)> + key_equals(const T &x) { + return [=](const std::pair<T, S> &ent) { + return ent.first == x; + }; + } + + template<typename T, typename S> + S + values(const std::pair<T, S> &ent) { + return ent.second; + } + + template<typename T> + std::function<bool (const T &)> + is_zero() { + return [](const T &x) { + return x == 0; + }; + } +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/compat.hpp b/src/gallium/state_trackers/clover/core/compat.hpp new file mode 100644 index 00000000000..c0057af3258 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/compat.hpp @@ -0,0 +1,290 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_COMPAT_HPP__ +#define __CORE_COMPAT_HPP__ + +#include <new> +#include <cstring> +#include <cstdlib> +#include <string> +#include <stdint.h> + + +namespace clover { + namespace compat { + // XXX - For cases where we can't rely on STL... I.e. the + // interface between code compiled as C++98 and C++11 + // source. Get rid of this as soon as everything can be + // compiled as C++11. + + template<typename T> + class vector { + protected: + static T * + alloc(int n, const T *q, int m) { + T *p = reinterpret_cast<T *>(std::malloc(n * sizeof(T))); + + for (int i = 0; i < m; ++i) + new(&p[i]) T(q[i]); + + return p; + } + + static void + free(int n, T *p) { + for (int i = 0; i < n; ++i) + p[i].~T(); + + std::free(p); + } + + public: + vector() : p(NULL), n(0) { + } + + vector(const vector &v) : p(alloc(v.n, v.p, v.n)), n(v.n) { + } + + vector(T *p, size_t n) : p(alloc(n, p, n)), n(n) { + } + + template<typename C> + vector(const C &v) : + p(alloc(v.size(), &*v.begin(), v.size())), n(v.size()) { + } + + ~vector() { + free(n, p); + } + + vector & + operator=(const vector &v) { + free(n, p); + + p = alloc(v.n, v.p, v.n); + n = v.n; + + return *this; + } + + void + reserve(size_t m) { + if (n < m) { + T *q = alloc(m, p, n); + free(n, p); + + p = q; + n = m; + } + } + + void + resize(size_t m, T x = T()) { + size_t n = size(); + + reserve(m); + + for (size_t i = n; i < m; ++i) + new(&p[i]) T(x); + } + + void + push_back(const T &x) { + size_t n = size(); + reserve(n + 1); + new(&p[n]) T(x); + } + + size_t + size() const { + return n; + } + + T * + begin() { + return p; + } + + const T * + begin() const { + return p; + } + + T * + end() { + return p + n; + } + + const T * + end() const { + return p + n; + } + + T & + operator[](int i) { + return p[i]; + } + + const T & + operator[](int i) const { + return p[i]; + } + + private: + T *p; + size_t n; + }; + + template<typename T> + class vector_ref { + public: + vector_ref(T *p, size_t n) : p(p), n(n) { + } + + template<typename C> + vector_ref(C &v) : p(&*v.begin()), n(v.size()) { + } + + size_t + size() const { + return n; + } + + T * + begin() { + return p; + } + + const T * + begin() const { + return p; + } + + T * + end() { + return p + n; + } + + const T * + end() const { + return p + n; + } + + T & + operator[](int i) { + return p[i]; + } + + const T & + operator[](int i) const { + return p[i]; + } + + private: + T *p; + size_t n; + }; + + class istream { + public: + typedef vector_ref<const unsigned char> buffer_t; + + class error { + public: + virtual ~error() {} + }; + + istream(const buffer_t &buf) : buf(buf), offset(0) {} + + void + read(char *p, size_t n) { + if (offset + n > buf.size()) + throw error(); + + std::memcpy(p, buf.begin() + offset, n); + offset += n; + } + + private: + const buffer_t &buf; + size_t offset; + }; + + class ostream { + public: + typedef vector<unsigned char> buffer_t; + + ostream(buffer_t &buf) : buf(buf), offset(buf.size()) {} + + void + write(const char *p, size_t n) { + buf.resize(offset + n); + std::memcpy(buf.begin() + offset, p, n); + offset += n; + } + + private: + buffer_t &buf; + size_t offset; + }; + + class string : public vector_ref<const char> { + public: + string(const char *p) : vector_ref(p, std::strlen(p)) { + } + + template<typename C> + string(const C &v) : vector_ref(v) { + } + + operator std::string() const { + return std::string(begin(), end()); + } + + const char * + find(const string &s) const { + for (size_t i = 0; i + s.size() < size(); ++i) { + if (!std::memcmp(begin() + i, s.begin(), s.size())) + return begin() + i; + } + + return end(); + } + }; + + template<typename T> + bool + operator==(const vector_ref<T> &a, const vector_ref<T> &b) { + if (a.size() != b.size()) + return false; + + for (size_t i = 0; i < a.size(); ++i) + if (a[i] != b[i]) + return false; + + return true; + } + } +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/compiler.hpp b/src/gallium/state_trackers/clover/core/compiler.hpp new file mode 100644 index 00000000000..a3998d5e2fb --- /dev/null +++ b/src/gallium/state_trackers/clover/core/compiler.hpp @@ -0,0 +1,53 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_COMPILER_HPP__ +#define __CORE_COMPILER_HPP__ + +#include "core/compat.hpp" +#include "core/module.hpp" + +namespace clover { + class build_error { + public: + build_error(const compat::string &log) : log(log) { + } + + virtual ~build_error() { + } + + compat::string what() { + return log; + } + + private: + compat::vector<char> log; + }; + + module compile_program_llvm(const compat::string &source, + const compat::string &target); + + module compile_program_tgsi(const compat::string &source, + const compat::string &target); +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/context.cpp b/src/gallium/state_trackers/clover/core/context.cpp new file mode 100644 index 00000000000..6e09a1acae0 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/context.cpp @@ -0,0 +1,37 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <algorithm> + +#include "core/context.hpp" + +using namespace clover; + +_cl_context::_cl_context(const std::vector<cl_context_properties> &props, + const std::vector<device *> &devs) : + devs(devs), __props(props) { +} + +bool +_cl_context::has_device(clover::device *dev) const { + return std::count(devs.begin(), devs.end(), dev); +} diff --git a/src/gallium/state_trackers/clover/core/context.hpp b/src/gallium/state_trackers/clover/core/context.hpp new file mode 100644 index 00000000000..d783fb6b14b --- /dev/null +++ b/src/gallium/state_trackers/clover/core/context.hpp @@ -0,0 +1,51 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_CONTEXT_HPP__ +#define __CORE_CONTEXT_HPP__ + +#include "core/base.hpp" +#include "core/device.hpp" + +namespace clover { + typedef struct _cl_context context; +} + +struct _cl_context : public clover::ref_counter { +public: + _cl_context(const std::vector<cl_context_properties> &props, + const std::vector<clover::device *> &devs); + _cl_context(const _cl_context &ctx) = delete; + + bool has_device(clover::device *dev) const; + + const std::vector<cl_context_properties> &props() const { + return __props; + } + + const std::vector<clover::device *> devs; + +private: + std::vector<cl_context_properties> __props; +}; + +#endif diff --git a/src/gallium/state_trackers/clover/core/device.cpp b/src/gallium/state_trackers/clover/core/device.cpp new file mode 100644 index 00000000000..8390f3f4abb --- /dev/null +++ b/src/gallium/state_trackers/clover/core/device.cpp @@ -0,0 +1,179 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/device.hpp" +#include "pipe/p_screen.h" +#include "pipe/p_state.h" + +using namespace clover; + +namespace { + template<typename T> + std::vector<T> + get_compute_param(pipe_screen *pipe, pipe_compute_cap cap) { + int sz = pipe->get_compute_param(pipe, cap, NULL); + std::vector<T> v(sz / sizeof(T)); + + pipe->get_compute_param(pipe, cap, &v.front()); + return v; + } +} + +_cl_device_id::_cl_device_id(pipe_loader_device *ldev) : ldev(ldev) { + pipe = pipe_loader_create_screen(ldev, PIPE_SEARCH_DIR); + if (!pipe || !pipe->get_param(pipe, PIPE_CAP_COMPUTE)) + throw error(CL_INVALID_DEVICE); +} + +_cl_device_id::_cl_device_id(_cl_device_id &&dev) : pipe(dev.pipe), ldev(dev.ldev) { + dev.ldev = NULL; + dev.pipe = NULL; +} + +_cl_device_id::~_cl_device_id() { + if (pipe) + pipe->destroy(pipe); + if (ldev) + pipe_loader_release(&ldev, 1); +} + +cl_device_type +_cl_device_id::type() const { + switch (ldev->type) { + case PIPE_LOADER_DEVICE_SOFTWARE: + return CL_DEVICE_TYPE_CPU; + case PIPE_LOADER_DEVICE_PCI: + return CL_DEVICE_TYPE_GPU; + default: + assert(0); + return 0; + } +} + +cl_uint +_cl_device_id::vendor_id() const { + switch (ldev->type) { + case PIPE_LOADER_DEVICE_SOFTWARE: + return 0; + case PIPE_LOADER_DEVICE_PCI: + return ldev->pci.vendor_id; + default: + assert(0); + return 0; + } +} + +size_t +_cl_device_id::max_images_read() const { + return PIPE_MAX_SHADER_RESOURCES; +} + +size_t +_cl_device_id::max_images_write() const { + return PIPE_MAX_SHADER_RESOURCES; +} + +cl_uint +_cl_device_id::max_image_levels_2d() const { + return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_2D_LEVELS); +} + +cl_uint +_cl_device_id::max_image_levels_3d() const { + return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_3D_LEVELS); +} + +cl_uint +_cl_device_id::max_samplers() const { + return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE, + PIPE_SHADER_CAP_MAX_TEXTURE_SAMPLERS); +} + +cl_ulong +_cl_device_id::max_mem_global() const { + return get_compute_param<uint64_t>(pipe, + PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE)[0]; +} + +cl_ulong +_cl_device_id::max_mem_local() const { + return get_compute_param<uint64_t>(pipe, + PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE)[0]; +} + +cl_ulong +_cl_device_id::max_mem_input() const { + return get_compute_param<uint64_t>(pipe, + PIPE_COMPUTE_CAP_MAX_INPUT_SIZE)[0]; +} + +cl_ulong +_cl_device_id::max_const_buffer_size() const { + return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE, + PIPE_SHADER_CAP_MAX_CONSTS) * 16; +} + +cl_uint +_cl_device_id::max_const_buffers() const { + return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE, + PIPE_SHADER_CAP_MAX_CONST_BUFFERS); +} + +std::vector<size_t> +_cl_device_id::max_block_size() const { + return get_compute_param<uint64_t>(pipe, PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE); +} + +std::string +_cl_device_id::device_name() const { + return pipe->get_name(pipe); +} + +std::string +_cl_device_id::vendor_name() const { + return pipe->get_vendor(pipe); +} + +std::string +_cl_device_id::ir_target() const { + switch (pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE, + PIPE_SHADER_CAP_PREFERRED_IR)) { + case PIPE_SHADER_IR_TGSI: + return "tgsi"; + default: + assert(0); + return ""; + } +} + +device_registry::device_registry() { + int n = pipe_loader_probe(NULL, 0); + std::vector<pipe_loader_device *> ldevs(n); + + pipe_loader_probe(&ldevs.front(), n); + + for (pipe_loader_device *ldev : ldevs) { + try { + devs.emplace_back(ldev); + } catch (error &) {} + } +} diff --git a/src/gallium/state_trackers/clover/core/device.hpp b/src/gallium/state_trackers/clover/core/device.hpp new file mode 100644 index 00000000000..8f284ba5e42 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/device.hpp @@ -0,0 +1,107 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_DEVICE_HPP__ +#define __CORE_DEVICE_HPP__ + +#include <set> +#include <vector> + +#include "core/base.hpp" +#include "core/format.hpp" +#include "pipe-loader/pipe_loader.h" + +namespace clover { + typedef struct _cl_device_id device; + class root_resource; + class hard_event; +} + +struct _cl_device_id { +public: + _cl_device_id(pipe_loader_device *ldev); + _cl_device_id(_cl_device_id &&dev); + _cl_device_id(const _cl_device_id &dev) = delete; + ~_cl_device_id(); + + cl_device_type type() const; + cl_uint vendor_id() const; + size_t max_images_read() const; + size_t max_images_write() const; + cl_uint max_image_levels_2d() const; + cl_uint max_image_levels_3d() const; + cl_uint max_samplers() const; + cl_ulong max_mem_global() const; + cl_ulong max_mem_local() const; + cl_ulong max_mem_input() const; + cl_ulong max_const_buffer_size() const; + cl_uint max_const_buffers() const; + + std::vector<size_t> max_block_size() const; + std::string device_name() const; + std::string vendor_name() const; + std::string ir_target() const; + + friend struct _cl_command_queue; + friend class clover::root_resource; + friend class clover::hard_event; + friend std::set<cl_image_format> + clover::supported_formats(cl_context, cl_mem_object_type); + +private: + pipe_screen *pipe; + pipe_loader_device *ldev; +}; + +namespace clover { + /// + /// Container of all the compute devices that are available in the + /// system. + /// + class device_registry { + public: + typedef std::vector<device>::iterator iterator; + + device_registry(); + + iterator begin() { + return devs.begin(); + } + + iterator end() { + return devs.end(); + } + + device &front() { + return devs.front(); + } + + device &back() { + return devs.back(); + } + + protected: + std::vector<device> devs; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/event.cpp b/src/gallium/state_trackers/clover/core/event.cpp new file mode 100644 index 00000000000..aa287e9a0c9 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/event.cpp @@ -0,0 +1,175 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/event.hpp" +#include "pipe/p_screen.h" + +using namespace clover; + +_cl_event::_cl_event(clover::context &ctx, + std::vector<clover::event *> deps, + action action_ok, action action_fail) : + ctx(ctx), __status(0), wait_count(1), + action_ok(action_ok), action_fail(action_fail) { + for (auto ev : deps) + ev->chain(this); +} + +_cl_event::~_cl_event() { +} + +void +_cl_event::trigger() { + if (!--wait_count) { + action_ok(*this); + + while (!__chain.empty()) { + __chain.back()->trigger(); + __chain.pop_back(); + } + } +} + +void +_cl_event::abort(cl_int status) { + __status = status; + action_fail(*this); + + while (!__chain.empty()) { + __chain.back()->abort(status); + __chain.pop_back(); + } +} + +bool +_cl_event::signalled() const { + return !wait_count; +} + +void +_cl_event::chain(clover::event *ev) { + if (wait_count) { + ev->wait_count++; + __chain.push_back(ev); + ev->deps.push_back(this); + } +} + +hard_event::hard_event(clover::command_queue &q, cl_command_type command, + std::vector<clover::event *> deps, action action) : + _cl_event(q.ctx, deps, action, [](event &ev){}), + __queue(q), __command(command), __fence(NULL) { + q.sequence(this); + trigger(); +} + +hard_event::~hard_event() { + pipe_screen *screen = queue()->dev.pipe; + screen->fence_reference(screen, &__fence, NULL); +} + +cl_int +hard_event::status() const { + pipe_screen *screen = queue()->dev.pipe; + + if (__status < 0) + return __status; + + else if (!__fence) + return CL_QUEUED; + + else if (!screen->fence_signalled(screen, __fence)) + return CL_SUBMITTED; + + else + return CL_COMPLETE; +} + +cl_command_queue +hard_event::queue() const { + return &__queue; +} + +cl_command_type +hard_event::command() const { + return __command; +} + +void +hard_event::wait() const { + pipe_screen *screen = queue()->dev.pipe; + + if (status() == CL_QUEUED) + queue()->flush(); + + if (!__fence || + !screen->fence_finish(screen, __fence, PIPE_TIMEOUT_INFINITE)) + throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST); +} + +void +hard_event::fence(pipe_fence_handle *fence) { + pipe_screen *screen = queue()->dev.pipe; + screen->fence_reference(screen, &__fence, fence); +} + +soft_event::soft_event(clover::context &ctx, + std::vector<clover::event *> deps, + bool __trigger, action action) : + _cl_event(ctx, deps, action, action) { + if (__trigger) + trigger(); +} + +cl_int +soft_event::status() const { + if (__status < 0) + return __status; + + else if (!signalled() || + any_of([](const ref_ptr<event> &ev) { + return ev->status() != CL_COMPLETE; + }, deps.begin(), deps.end())) + return CL_SUBMITTED; + + else + return CL_COMPLETE; +} + +cl_command_queue +soft_event::queue() const { + return NULL; +} + +cl_command_type +soft_event::command() const { + return CL_COMMAND_USER; +} + +void +soft_event::wait() const { + for (auto ev : deps) + ev->wait(); + + if (status() != CL_COMPLETE) + throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST); +} diff --git a/src/gallium/state_trackers/clover/core/event.hpp b/src/gallium/state_trackers/clover/core/event.hpp new file mode 100644 index 00000000000..ea4ac4ae43c --- /dev/null +++ b/src/gallium/state_trackers/clover/core/event.hpp @@ -0,0 +1,138 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_EVENT_HPP__ +#define __CORE_EVENT_HPP__ + +#include <functional> + +#include "core/base.hpp" +#include "core/queue.hpp" + +namespace clover { + typedef struct _cl_event event; +} + +/// +/// Class that represents a task that might be executed asynchronously +/// at some point in the future. +/// +/// An event consists of a list of dependencies, a boolean signalled() +/// flag, and an associated task. An event is considered signalled as +/// soon as all its dependencies (if any) are signalled as well, and +/// the trigger() method is called; at that point the associated task +/// will be started through the specified \a action_ok. If the +/// abort() method is called instead, the specified \a action_fail is +/// executed and the associated task will never be started. Dependent +/// events will be aborted recursively. +/// +/// The execution status of the associated task can be queried using +/// the status() method, and it can be waited for completion using the +/// wait() method. +/// +struct _cl_event : public clover::ref_counter { +public: + typedef std::function<void (clover::event &)> action; + + _cl_event(clover::context &ctx, std::vector<clover::event *> deps, + action action_ok, action action_fail); + virtual ~_cl_event(); + + void trigger(); + void abort(cl_int status); + bool signalled() const; + + virtual cl_int status() const = 0; + virtual cl_command_queue queue() const = 0; + virtual cl_command_type command() const = 0; + virtual void wait() const = 0; + + clover::context &ctx; + +protected: + void chain(clover::event *ev); + + cl_int __status; + std::vector<clover::ref_ptr<clover::event>> deps; + +private: + unsigned wait_count; + action action_ok; + action action_fail; + std::vector<clover::ref_ptr<clover::event>> __chain; +}; + +namespace clover { + /// + /// Class that represents a task executed by a command queue. + /// + /// Similar to a normal clover::event. In addition it's associated + /// with a given command queue \a q and a given OpenCL \a command. + /// hard_event instances created for the same queue are implicitly + /// ordered with respect to each other, and they are implicitly + /// triggered on construction. + /// + /// A hard_event is considered complete when the associated + /// hardware task finishes execution. + /// + class hard_event : public event { + public: + hard_event(clover::command_queue &q, cl_command_type command, + std::vector<clover::event *> deps, + action action = [](event &){}); + ~hard_event(); + + virtual cl_int status() const; + virtual cl_command_queue queue() const; + virtual cl_command_type command() const; + virtual void wait() const; + + friend class ::_cl_command_queue; + + private: + virtual void fence(pipe_fence_handle *fence); + + clover::command_queue &__queue; + cl_command_type __command; + pipe_fence_handle *__fence; + }; + + /// + /// Class that represents a software event. + /// + /// A soft_event is not associated with any specific hardware task + /// or command queue. It's considered complete as soon as all its + /// dependencies finish execution. + /// + class soft_event : public event { + public: + soft_event(clover::context &ctx, std::vector<clover::event *> deps, + bool trigger, action action = [](event &){}); + + virtual cl_int status() const; + virtual cl_command_queue queue() const; + virtual cl_command_type command() const; + virtual void wait() const; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/format.cpp b/src/gallium/state_trackers/clover/core/format.cpp new file mode 100644 index 00000000000..8f6e14d6567 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/format.cpp @@ -0,0 +1,167 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <algorithm> + +#include "core/format.hpp" +#include "core/memory.hpp" +#include "pipe/p_screen.h" +#include "pipe/p_context.h" + +namespace clover { + static const std::map<cl_image_format, pipe_format> formats { + { { CL_BGRA, CL_UNORM_INT8 }, PIPE_FORMAT_B8G8R8A8_UNORM }, + { { CL_ARGB, CL_UNORM_INT8 }, PIPE_FORMAT_A8R8G8B8_UNORM }, + { { CL_RGB, CL_UNORM_SHORT_565 }, PIPE_FORMAT_B5G6R5_UNORM }, + { { CL_LUMINANCE, CL_UNORM_INT8 }, PIPE_FORMAT_L8_UNORM }, + { { CL_A, CL_UNORM_INT8 }, PIPE_FORMAT_A8_UNORM }, + { { CL_INTENSITY, CL_UNORM_INT8 }, PIPE_FORMAT_I8_UNORM }, + { { CL_LUMINANCE, CL_UNORM_INT16 }, PIPE_FORMAT_L16_UNORM }, + { { CL_R, CL_FLOAT }, PIPE_FORMAT_R32_FLOAT }, + { { CL_RG, CL_FLOAT }, PIPE_FORMAT_R32G32_FLOAT }, + { { CL_RGB, CL_FLOAT }, PIPE_FORMAT_R32G32B32_FLOAT }, + { { CL_RGBA, CL_FLOAT }, PIPE_FORMAT_R32G32B32A32_FLOAT }, + { { CL_R, CL_UNORM_INT16 }, PIPE_FORMAT_R16_UNORM }, + { { CL_RG, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16_UNORM }, + { { CL_RGB, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16B16_UNORM }, + { { CL_RGBA, CL_UNORM_INT16 }, PIPE_FORMAT_R16G16B16A16_UNORM }, + { { CL_R, CL_SNORM_INT16 }, PIPE_FORMAT_R16_SNORM }, + { { CL_RG, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16_SNORM }, + { { CL_RGB, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16B16_SNORM }, + { { CL_RGBA, CL_SNORM_INT16 }, PIPE_FORMAT_R16G16B16A16_SNORM }, + { { CL_R, CL_UNORM_INT8 }, PIPE_FORMAT_R8_UNORM }, + { { CL_RG, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8_UNORM }, + { { CL_RGB, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8_UNORM }, + { { CL_RGBA, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8A8_UNORM }, + { { CL_R, CL_SNORM_INT8 }, PIPE_FORMAT_R8_SNORM }, + { { CL_RG, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8_SNORM }, + { { CL_RGB, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8B8_SNORM }, + { { CL_RGBA, CL_SNORM_INT8 }, PIPE_FORMAT_R8G8B8A8_SNORM }, + { { CL_R, CL_HALF_FLOAT }, PIPE_FORMAT_R16_FLOAT }, + { { CL_RG, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16_FLOAT }, + { { CL_RGB, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16B16_FLOAT }, + { { CL_RGBA, CL_HALF_FLOAT }, PIPE_FORMAT_R16G16B16A16_FLOAT }, + { { CL_RGBx, CL_UNORM_SHORT_555 }, PIPE_FORMAT_B5G5R5X1_UNORM }, + { { CL_RGBx, CL_UNORM_INT8 }, PIPE_FORMAT_R8G8B8X8_UNORM }, + { { CL_A, CL_UNORM_INT16 }, PIPE_FORMAT_A16_UNORM }, + { { CL_INTENSITY, CL_UNORM_INT16 }, PIPE_FORMAT_I16_UNORM }, + { { CL_LUMINANCE, CL_SNORM_INT8 }, PIPE_FORMAT_L8_SNORM }, + { { CL_INTENSITY, CL_SNORM_INT8 }, PIPE_FORMAT_I8_SNORM }, + { { CL_A, CL_SNORM_INT16 }, PIPE_FORMAT_A16_SNORM }, + { { CL_LUMINANCE, CL_SNORM_INT16 }, PIPE_FORMAT_L16_SNORM }, + { { CL_INTENSITY, CL_SNORM_INT16 }, PIPE_FORMAT_I16_SNORM }, + { { CL_A, CL_HALF_FLOAT }, PIPE_FORMAT_A16_FLOAT }, + { { CL_LUMINANCE, CL_HALF_FLOAT }, PIPE_FORMAT_L16_FLOAT }, + { { CL_INTENSITY, CL_HALF_FLOAT }, PIPE_FORMAT_I16_FLOAT }, + { { CL_A, CL_FLOAT }, PIPE_FORMAT_A32_FLOAT }, + { { CL_LUMINANCE, CL_FLOAT }, PIPE_FORMAT_L32_FLOAT }, + { { CL_INTENSITY, CL_FLOAT }, PIPE_FORMAT_I32_FLOAT }, + { { CL_RA, CL_UNORM_INT8 }, PIPE_FORMAT_R8A8_UNORM }, + { { CL_R, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8_UINT }, + { { CL_RG, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8_UINT }, + { { CL_RGB, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8B8_UINT }, + { { CL_RGBA, CL_UNSIGNED_INT8 }, PIPE_FORMAT_R8G8B8A8_UINT }, + { { CL_R, CL_SIGNED_INT8 }, PIPE_FORMAT_R8_SINT }, + { { CL_RG, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8_SINT }, + { { CL_RGB, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8B8_SINT }, + { { CL_RGBA, CL_SIGNED_INT8 }, PIPE_FORMAT_R8G8B8A8_SINT }, + { { CL_R, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16_UINT }, + { { CL_RG, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16_UINT }, + { { CL_RGB, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16B16_UINT }, + { { CL_RGBA, CL_UNSIGNED_INT16 }, PIPE_FORMAT_R16G16B16A16_UINT }, + { { CL_R, CL_SIGNED_INT16 }, PIPE_FORMAT_R16_SINT }, + { { CL_RG, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16_SINT }, + { { CL_RGB, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16B16_SINT }, + { { CL_RGBA, CL_SIGNED_INT16 }, PIPE_FORMAT_R16G16B16A16_SINT }, + { { CL_R, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32_UINT }, + { { CL_RG, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32_UINT }, + { { CL_RGB, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32B32_UINT }, + { { CL_RGBA, CL_UNSIGNED_INT32 }, PIPE_FORMAT_R32G32B32A32_UINT }, + { { CL_R, CL_SIGNED_INT32 }, PIPE_FORMAT_R32_SINT }, + { { CL_RG, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32_SINT }, + { { CL_RGB, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32B32_SINT }, + { { CL_RGBA, CL_SIGNED_INT32 }, PIPE_FORMAT_R32G32B32A32_SINT }, + { { CL_A, CL_UNSIGNED_INT8 }, PIPE_FORMAT_A8_UINT }, + { { CL_INTENSITY, CL_UNSIGNED_INT8 }, PIPE_FORMAT_I8_UINT }, + { { CL_LUMINANCE, CL_UNSIGNED_INT8 }, PIPE_FORMAT_L8_UINT }, + { { CL_A, CL_SIGNED_INT8 }, PIPE_FORMAT_A8_SINT }, + { { CL_INTENSITY, CL_SIGNED_INT8 }, PIPE_FORMAT_I8_SINT }, + { { CL_LUMINANCE, CL_SIGNED_INT8 }, PIPE_FORMAT_L8_SINT }, + { { CL_A, CL_UNSIGNED_INT16 }, PIPE_FORMAT_A16_UINT }, + { { CL_INTENSITY, CL_UNSIGNED_INT16 }, PIPE_FORMAT_I16_UINT }, + { { CL_LUMINANCE, CL_UNSIGNED_INT16 }, PIPE_FORMAT_L16_UINT }, + { { CL_A, CL_SIGNED_INT16 }, PIPE_FORMAT_A16_SINT }, + { { CL_INTENSITY, CL_SIGNED_INT16 }, PIPE_FORMAT_I16_SINT }, + { { CL_LUMINANCE, CL_SIGNED_INT16 }, PIPE_FORMAT_L16_SINT }, + { { CL_A, CL_UNSIGNED_INT32 }, PIPE_FORMAT_A32_UINT }, + { { CL_INTENSITY, CL_UNSIGNED_INT32 }, PIPE_FORMAT_I32_UINT }, + { { CL_LUMINANCE, CL_UNSIGNED_INT32 }, PIPE_FORMAT_L32_UINT }, + { { CL_A, CL_SIGNED_INT32 }, PIPE_FORMAT_A32_SINT }, + { { CL_INTENSITY, CL_SIGNED_INT32 }, PIPE_FORMAT_I32_SINT }, + { { CL_LUMINANCE, CL_SIGNED_INT32 }, PIPE_FORMAT_L32_SINT } + }; + + pipe_texture_target + translate_target(cl_mem_object_type type) { + switch (type) { + case CL_MEM_OBJECT_BUFFER: + return PIPE_BUFFER; + case CL_MEM_OBJECT_IMAGE2D: + return PIPE_TEXTURE_2D; + case CL_MEM_OBJECT_IMAGE3D: + return PIPE_TEXTURE_3D; + default: + throw error(CL_INVALID_VALUE); + } + } + + pipe_format + translate_format(const cl_image_format &format) { + auto it = formats.find(format); + + if (it == formats.end()) + throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED); + + return it->second; + } + + std::set<cl_image_format> + supported_formats(cl_context ctx, cl_mem_object_type type) { + std::set<cl_image_format> s; + pipe_texture_target target = translate_target(type); + unsigned bindings = (PIPE_BIND_SAMPLER_VIEW | + PIPE_BIND_COMPUTE_RESOURCE | + PIPE_BIND_TRANSFER_READ | + PIPE_BIND_TRANSFER_WRITE); + + for (auto f : formats) { + if (std::all_of(ctx->devs.begin(), ctx->devs.end(), + [=](const device *dev) { + return dev->pipe->is_format_supported( + dev->pipe, f.second, target, 1, bindings); + })) + s.insert(f.first); + } + + return s; + } +} diff --git a/src/gallium/state_trackers/clover/core/format.hpp b/src/gallium/state_trackers/clover/core/format.hpp new file mode 100644 index 00000000000..a24cbf37621 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/format.hpp @@ -0,0 +1,51 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_FORMAT_HPP__ +#define __CORE_FORMAT_HPP__ + +#include <set> + +#include "core/base.hpp" +#include "pipe/p_defines.h" +#include "pipe/p_format.h" + +namespace clover { + pipe_texture_target translate_target(cl_mem_object_type type); + pipe_format translate_format(const cl_image_format &format); + + /// + /// Return all the image formats supported by a given context for + /// the given memory object type. + /// + std::set<cl_image_format> supported_formats(cl_context ctx, + cl_mem_object_type type); +} + +static inline bool +operator<(const cl_image_format &a, const cl_image_format &b) { + return (a.image_channel_order != b.image_channel_order ? + a.image_channel_order < b.image_channel_order : + a.image_channel_data_type < b.image_channel_data_type); +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/geometry.hpp b/src/gallium/state_trackers/clover/core/geometry.hpp new file mode 100644 index 00000000000..027264e72f0 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/geometry.hpp @@ -0,0 +1,72 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_GEOMETRY_HPP__ +#define __CORE_GEOMETRY_HPP__ + +#include <array> +#include <algorithm> + +namespace clover { + /// + /// N-dimensional coordinate array. + /// + template<typename T, int N> + class point { + public: + point() : a() { + } + + point(std::initializer_list<T> v) { + auto it = std::copy(v.begin(), v.end(), a.begin()); + std::fill(it, a.end(), 0); + } + + point(const T *v) { + std::copy(v, v + N, a.begin()); + } + + T &operator[](int i) { + return a[i]; + } + + const T &operator[](int i) const { + return a[i]; + } + + point operator+(const point &p) const { + point q; + std::transform(a.begin(), a.end(), p.a.begin(), + q.a.begin(), std::plus<T>()); + return q; + } + + T operator()(const point &p) const { + return std::inner_product(p.a.begin(), p.a.end(), a.begin(), 0); + } + + protected: + std::array<T, N> a; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/kernel.cpp b/src/gallium/state_trackers/clover/core/kernel.cpp new file mode 100644 index 00000000000..6fa8bd63453 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/kernel.cpp @@ -0,0 +1,393 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/kernel.hpp" +#include "core/resource.hpp" +#include "pipe/p_context.h" + +using namespace clover; + +_cl_kernel::_cl_kernel(clover::program &prog, + const std::string &name, + const std::vector<clover::module::argument> &args) : + prog(prog), __name(name), exec(*this) { + for (auto arg : args) { + if (arg.type == module::argument::scalar) + this->args.emplace_back(new scalar_argument(arg.size)); + else if (arg.type == module::argument::global) + this->args.emplace_back(new global_argument(arg.size)); + else if (arg.type == module::argument::local) + this->args.emplace_back(new local_argument()); + else if (arg.type == module::argument::constant) + this->args.emplace_back(new constant_argument()); + else if (arg.type == module::argument::image2d_rd || + arg.type == module::argument::image3d_rd) + this->args.emplace_back(new image_rd_argument()); + else if (arg.type == module::argument::image2d_wr || + arg.type == module::argument::image3d_wr) + this->args.emplace_back(new image_wr_argument()); + else if (arg.type == module::argument::sampler) + this->args.emplace_back(new sampler_argument()); + else + throw error(CL_INVALID_KERNEL_DEFINITION); + } +} + +template<typename T, typename V> +static inline std::vector<T> +pad_vector(clover::command_queue &q, const V &v, T x) { + std::vector<T> w { v.begin(), v.end() }; + w.resize(q.dev.max_block_size().size(), x); + return w; +} + +void +_cl_kernel::launch(clover::command_queue &q, + const std::vector<size_t> &grid_offset, + const std::vector<size_t> &grid_size, + const std::vector<size_t> &block_size) { + void *st = exec.bind(&q); + auto g_handles = map([&](size_t h) { return (uint32_t *)&exec.input[h]; }, + exec.g_handles.begin(), exec.g_handles.end()); + + q.pipe->bind_compute_state(q.pipe, st); + q.pipe->bind_compute_sampler_states(q.pipe, 0, exec.samplers.size(), + exec.samplers.data()); + q.pipe->set_compute_sampler_views(q.pipe, 0, exec.sviews.size(), + exec.sviews.data()); + q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(), + exec.resources.data()); + q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(), + exec.g_buffers.data(), g_handles.data()); + + q.pipe->launch_grid(q.pipe, + pad_vector<uint>(q, block_size, 1).data(), + pad_vector<uint>(q, grid_size, 1).data(), + module(q).sym(__name).offset, + exec.input.data()); + + q.pipe->set_global_binding(q.pipe, 0, exec.g_buffers.size(), NULL, NULL); + q.pipe->set_compute_resources(q.pipe, 0, exec.resources.size(), NULL); + q.pipe->set_compute_sampler_views(q.pipe, 0, exec.sviews.size(), NULL); + q.pipe->bind_compute_sampler_states(q.pipe, 0, exec.samplers.size(), NULL); + exec.unbind(); +} + +size_t +_cl_kernel::mem_local() const { + size_t sz = 0; + + for (auto &arg : args) { + if (dynamic_cast<local_argument *>(arg.get())) + sz += arg->storage(); + } + + return sz; +} + +size_t +_cl_kernel::mem_private() const { + return 0; +} + +size_t +_cl_kernel::max_block_size() const { + return SIZE_MAX; +} + +const std::string & +_cl_kernel::name() const { + return __name; +} + +std::vector<size_t> +_cl_kernel::block_size() const { + return { 0, 0, 0 }; +} + +const clover::module & +_cl_kernel::module(const clover::command_queue &q) const { + return prog.binaries().find(&q.dev)->second; +} + + +_cl_kernel::exec_context::exec_context(clover::kernel &kern) : + kern(kern), q(NULL), mem_local(0), st(NULL) { +} + +_cl_kernel::exec_context::~exec_context() { + if (st) + q->pipe->delete_compute_state(q->pipe, st); +} + +void * +_cl_kernel::exec_context::bind(clover::command_queue *__q) { + std::swap(q, __q); + + for (auto &arg : kern.args) + arg->bind(*this); + + // Create a new compute state if anything changed. + if (!st || q != __q || + cs.req_local_mem != mem_local || + cs.req_input_mem != input.size()) { + if (st) + __q->pipe->delete_compute_state(__q->pipe, st); + + cs.prog = kern.module(*q).sec(module::section::text).data.begin(); + cs.req_local_mem = mem_local; + cs.req_input_mem = input.size(); + st = q->pipe->create_compute_state(q->pipe, &cs); + } + + return st; +} + +void +_cl_kernel::exec_context::unbind() { + for (auto &arg : kern.args) + arg->unbind(*this); + + input.clear(); + samplers.clear(); + sviews.clear(); + resources.clear(); + g_buffers.clear(); + g_handles.clear(); + mem_local = 0; +} + +_cl_kernel::argument::argument(size_t size) : + __size(size), __set(false) { +} + +bool +_cl_kernel::argument::set() const { + return __set; +} + +size_t +_cl_kernel::argument::storage() const { + return 0; +} + +_cl_kernel::scalar_argument::scalar_argument(size_t size) : + argument(size) { +} + +void +_cl_kernel::scalar_argument::set(size_t size, const void *value) { + if (size != __size) + throw error(CL_INVALID_ARG_SIZE); + + v = { (uint8_t *)value, (uint8_t *)value + size }; + __set = true; +} + +void +_cl_kernel::scalar_argument::bind(exec_context &ctx) { + ctx.input.insert(ctx.input.end(), v.begin(), v.end()); +} + +void +_cl_kernel::scalar_argument::unbind(exec_context &ctx) { +} + +_cl_kernel::global_argument::global_argument(size_t size) : + argument(size) { +} + +void +_cl_kernel::global_argument::set(size_t size, const void *value) { + if (size != sizeof(cl_mem)) + throw error(CL_INVALID_ARG_SIZE); + + obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value); + __set = true; +} + +void +_cl_kernel::global_argument::bind(exec_context &ctx) { + size_t offset = ctx.input.size(); + size_t idx = ctx.g_buffers.size(); + + ctx.input.resize(offset + __size); + + ctx.g_buffers.resize(idx + 1); + ctx.g_buffers[idx] = obj->resource(ctx.q).pipe; + + ctx.g_handles.resize(idx + 1); + ctx.g_handles[idx] = offset; +} + +void +_cl_kernel::global_argument::unbind(exec_context &ctx) { +} + +_cl_kernel::local_argument::local_argument() : + argument(sizeof(uint32_t)) { +} + +size_t +_cl_kernel::local_argument::storage() const { + return __storage; +} + +void +_cl_kernel::local_argument::set(size_t size, const void *value) { + if (value) + throw error(CL_INVALID_ARG_VALUE); + + __storage = size; + __set = true; +} + +void +_cl_kernel::local_argument::bind(exec_context &ctx) { + size_t offset = ctx.input.size(); + size_t ptr = ctx.mem_local; + + ctx.input.resize(offset + sizeof(uint32_t)); + *(uint32_t *)&ctx.input[offset] = ptr; + + ctx.mem_local += __storage; +} + +void +_cl_kernel::local_argument::unbind(exec_context &ctx) { +} + +_cl_kernel::constant_argument::constant_argument() : + argument(sizeof(uint32_t)) { +} + +void +_cl_kernel::constant_argument::set(size_t size, const void *value) { + if (size != sizeof(cl_mem)) + throw error(CL_INVALID_ARG_SIZE); + + obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value); + __set = true; +} + +void +_cl_kernel::constant_argument::bind(exec_context &ctx) { + size_t offset = ctx.input.size(); + size_t idx = ctx.resources.size(); + + ctx.input.resize(offset + sizeof(uint32_t)); + *(uint32_t *)&ctx.input[offset] = idx << 24; + + ctx.resources.resize(idx + 1); + ctx.resources[idx] = st = obj->resource(ctx.q).bind_surface(*ctx.q, false); +} + +void +_cl_kernel::constant_argument::unbind(exec_context &ctx) { + obj->resource(ctx.q).unbind_surface(*ctx.q, st); +} + +_cl_kernel::image_rd_argument::image_rd_argument() : + argument(sizeof(uint32_t)) { +} + +void +_cl_kernel::image_rd_argument::set(size_t size, const void *value) { + if (size != sizeof(cl_mem)) + throw error(CL_INVALID_ARG_SIZE); + + obj = dynamic_cast<clover::image *>(*(cl_mem *)value); + __set = true; +} + +void +_cl_kernel::image_rd_argument::bind(exec_context &ctx) { + size_t offset = ctx.input.size(); + size_t idx = ctx.sviews.size(); + + ctx.input.resize(offset + sizeof(uint32_t)); + *(uint32_t *)&ctx.input[offset] = idx; + + ctx.sviews.resize(idx + 1); + ctx.sviews[idx] = st = obj->resource(ctx.q).bind_sampler_view(*ctx.q); +} + +void +_cl_kernel::image_rd_argument::unbind(exec_context &ctx) { + obj->resource(ctx.q).unbind_sampler_view(*ctx.q, st); +} + +_cl_kernel::image_wr_argument::image_wr_argument() : + argument(sizeof(uint32_t)) { +} + +void +_cl_kernel::image_wr_argument::set(size_t size, const void *value) { + if (size != sizeof(cl_mem)) + throw error(CL_INVALID_ARG_SIZE); + + obj = dynamic_cast<clover::image *>(*(cl_mem *)value); + __set = true; +} + +void +_cl_kernel::image_wr_argument::bind(exec_context &ctx) { + size_t offset = ctx.input.size(); + size_t idx = ctx.resources.size(); + + ctx.input.resize(offset + sizeof(uint32_t)); + *(uint32_t *)&ctx.input[offset] = idx; + + ctx.resources.resize(idx + 1); + ctx.resources[idx] = st = obj->resource(ctx.q).bind_surface(*ctx.q, true); +} + +void +_cl_kernel::image_wr_argument::unbind(exec_context &ctx) { + obj->resource(ctx.q).unbind_surface(*ctx.q, st); +} + +_cl_kernel::sampler_argument::sampler_argument() : + argument(0) { +} + +void +_cl_kernel::sampler_argument::set(size_t size, const void *value) { + if (size != sizeof(cl_sampler)) + throw error(CL_INVALID_ARG_SIZE); + + obj = *(cl_sampler *)value; + __set = true; +} + +void +_cl_kernel::sampler_argument::bind(exec_context &ctx) { + size_t idx = ctx.samplers.size(); + + ctx.samplers.resize(idx + 1); + ctx.samplers[idx] = st = obj->bind(*ctx.q); +} + +void +_cl_kernel::sampler_argument::unbind(exec_context &ctx) { + obj->unbind(*ctx.q, st); +} diff --git a/src/gallium/state_trackers/clover/core/kernel.hpp b/src/gallium/state_trackers/clover/core/kernel.hpp new file mode 100644 index 00000000000..bc21de8094f --- /dev/null +++ b/src/gallium/state_trackers/clover/core/kernel.hpp @@ -0,0 +1,214 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_KERNEL_HPP__ +#define __CORE_KERNEL_HPP__ + +#include <memory> + +#include "core/base.hpp" +#include "core/program.hpp" +#include "core/memory.hpp" +#include "core/sampler.hpp" +#include "pipe/p_state.h" + +namespace clover { + typedef struct _cl_kernel kernel; + class argument; +} + +struct _cl_kernel : public clover::ref_counter { +private: + /// + /// Class containing all the state required to execute a compute + /// kernel. + /// + struct exec_context { + exec_context(clover::kernel &kern); + ~exec_context(); + + void *bind(clover::command_queue *q); + void unbind(); + + clover::kernel &kern; + clover::command_queue *q; + + std::vector<uint8_t> input; + std::vector<void *> samplers; + std::vector<pipe_sampler_view *> sviews; + std::vector<pipe_surface *> resources; + std::vector<pipe_resource *> g_buffers; + std::vector<size_t> g_handles; + size_t mem_local; + + private: + void *st; + pipe_compute_state cs; + }; + +public: + class argument { + public: + argument(size_t size); + + /// \a true if the argument has been set. + bool set() const; + + /// Argument size in the input buffer. + size_t size() const; + + /// Storage space required for the referenced object. + virtual size_t storage() const; + + /// Set this argument to some object. + virtual void set(size_t size, const void *value) = 0; + + /// Allocate the necessary resources to bind the specified + /// object to this argument, and update \a ctx accordingly. + virtual void bind(exec_context &ctx) = 0; + + /// Free any resources that were allocated in bind(). + virtual void unbind(exec_context &ctx) = 0; + + protected: + size_t __size; + bool __set; + }; + + _cl_kernel(clover::program &prog, + const std::string &name, + const std::vector<clover::module::argument> &args); + + void launch(clover::command_queue &q, + const std::vector<size_t> &grid_offset, + const std::vector<size_t> &grid_size, + const std::vector<size_t> &block_size); + + size_t mem_local() const; + size_t mem_private() const; + size_t max_block_size() const; + + const std::string &name() const; + std::vector<size_t> block_size() const; + + clover::program &prog; + std::vector<std::unique_ptr<argument>> args; + +private: + const clover::module & + module(const clover::command_queue &q) const; + + class scalar_argument : public argument { + public: + scalar_argument(size_t size); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + std::vector<uint8_t> v; + }; + + class global_argument : public argument { + public: + global_argument(size_t size); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + clover::buffer *obj; + }; + + class local_argument : public argument { + public: + local_argument(); + + virtual size_t storage() const; + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + size_t __storage; + }; + + class constant_argument : public argument { + public: + constant_argument(); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + clover::buffer *obj; + pipe_surface *st; + }; + + class image_rd_argument : public argument { + public: + image_rd_argument(); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + clover::image *obj; + pipe_sampler_view *st; + }; + + class image_wr_argument : public argument { + public: + image_wr_argument(); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + clover::image *obj; + pipe_surface *st; + }; + + class sampler_argument : public argument { + public: + sampler_argument(); + + virtual void set(size_t size, const void *value); + virtual void bind(exec_context &ctx); + virtual void unbind(exec_context &ctx); + + private: + clover::sampler *obj; + void *st; + }; + + std::string __name; + exec_context exec; +}; + +#endif diff --git a/src/gallium/state_trackers/clover/core/memory.cpp b/src/gallium/state_trackers/clover/core/memory.cpp new file mode 100644 index 00000000000..1bf12e3c36e --- /dev/null +++ b/src/gallium/state_trackers/clover/core/memory.cpp @@ -0,0 +1,198 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/memory.hpp" +#include "core/resource.hpp" + +using namespace clover; + +_cl_mem::_cl_mem(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr) : + ctx(ctx), __flags(flags), + __size(size), __host_ptr(host_ptr), + __destroy_notify([]{}), + data((char *)host_ptr, (host_ptr ? size : 0)) { +} + +_cl_mem::~_cl_mem() { + __destroy_notify(); +} + +void +_cl_mem::destroy_notify(std::function<void ()> f) { + __destroy_notify = f; +} + +cl_mem_flags +_cl_mem::flags() const { + return __flags; +} + +size_t +_cl_mem::size() const { + return __size; +} + +void * +_cl_mem::host_ptr() const { + return __host_ptr; +} + +buffer::buffer(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr) : + memory_obj(ctx, flags, size, host_ptr) { +} + +cl_mem_object_type +buffer::type() const { + return CL_MEM_OBJECT_BUFFER; +} + +root_buffer::root_buffer(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr) : + buffer(ctx, flags, size, host_ptr) { +} + +clover::resource & +root_buffer::resource(cl_command_queue q) { + // Create a new resource if there's none for this device yet. + if (!resources.count(&q->dev)) { + auto r = (!resources.empty() ? + new root_resource(q->dev, *this, *resources.begin()->second) : + new root_resource(q->dev, *this, data)); + + resources.insert(std::make_pair(&q->dev, + std::unique_ptr<root_resource>(r))); + data.clear(); + } + + return *resources.find(&q->dev)->second; +} + +sub_buffer::sub_buffer(clover::root_buffer &parent, cl_mem_flags flags, + size_t offset, size_t size) : + buffer(parent.ctx, flags, size, + (char *)parent.host_ptr() + offset), + parent(parent), __offset(offset) { +} + +clover::resource & +sub_buffer::resource(cl_command_queue q) { + // Create a new resource if there's none for this device yet. + if (!resources.count(&q->dev)) { + auto r = new sub_resource(parent.resource(q), { offset() }); + + resources.insert(std::make_pair(&q->dev, + std::unique_ptr<sub_resource>(r))); + } + + return *resources.find(&q->dev)->second; +} + +size_t +sub_buffer::offset() const { + return __offset; +} + +image::image(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, size_t size, + void *host_ptr) : + memory_obj(ctx, flags, size, host_ptr), + __format(*format), __width(width), __height(height), __depth(depth), + __row_pitch(row_pitch), __slice_pitch(slice_pitch) { +} + +clover::resource & +image::resource(cl_command_queue q) { + // Create a new resource if there's none for this device yet. + if (!resources.count(&q->dev)) { + auto r = (!resources.empty() ? + new root_resource(q->dev, *this, *resources.begin()->second) : + new root_resource(q->dev, *this, data)); + + resources.insert(std::make_pair(&q->dev, + std::unique_ptr<root_resource>(r))); + data.clear(); + } + + return *resources.find(&q->dev)->second; +} + +cl_image_format +image::format() const { + return __format; +} + +size_t +image::width() const { + return __width; +} + +size_t +image::height() const { + return __height; +} + +size_t +image::depth() const { + return __depth; +} + +size_t +image::row_pitch() const { + return __row_pitch; +} + +size_t +image::slice_pitch() const { + return __slice_pitch; +} + +image2d::image2d(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, size_t width, + size_t height, size_t row_pitch, + void *host_ptr) : + image(ctx, flags, format, width, height, 0, + row_pitch, 0, height * row_pitch, host_ptr) { +} + +cl_mem_object_type +image2d::type() const { + return CL_MEM_OBJECT_IMAGE2D; +} + +image3d::image3d(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, + void *host_ptr) : + image(ctx, flags, format, width, height, depth, + row_pitch, slice_pitch, depth * slice_pitch, + host_ptr) { +} + +cl_mem_object_type +image3d::type() const { + return CL_MEM_OBJECT_IMAGE3D; +} diff --git a/src/gallium/state_trackers/clover/core/memory.hpp b/src/gallium/state_trackers/clover/core/memory.hpp new file mode 100644 index 00000000000..96f70e931bc --- /dev/null +++ b/src/gallium/state_trackers/clover/core/memory.hpp @@ -0,0 +1,157 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_MEMORY_HPP__ +#define __CORE_MEMORY_HPP__ + +#include <functional> +#include <map> +#include <memory> + +#include "core/base.hpp" +#include "core/queue.hpp" + +namespace clover { + typedef struct _cl_mem memory_obj; + + class resource; + class sub_resource; +} + +struct _cl_mem : public clover::ref_counter { +protected: + _cl_mem(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr); + _cl_mem(const _cl_mem &obj) = delete; + +public: + virtual ~_cl_mem(); + + virtual cl_mem_object_type type() const = 0; + virtual clover::resource &resource(cl_command_queue q) = 0; + + void destroy_notify(std::function<void ()> f); + cl_mem_flags flags() const; + size_t size() const; + void *host_ptr() const; + + clover::context &ctx; + +private: + cl_mem_flags __flags; + size_t __size; + void *__host_ptr; + std::function<void ()> __destroy_notify; + +protected: + std::string data; +}; + +namespace clover { + struct buffer : public memory_obj { + protected: + buffer(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr); + + public: + virtual cl_mem_object_type type() const; + }; + + struct root_buffer : public buffer { + public: + root_buffer(clover::context &ctx, cl_mem_flags flags, + size_t size, void *host_ptr); + + virtual clover::resource &resource(cl_command_queue q); + + private: + std::map<clover::device *, + std::unique_ptr<clover::root_resource>> resources; + }; + + struct sub_buffer : public buffer { + public: + sub_buffer(clover::root_buffer &parent, cl_mem_flags flags, + size_t offset, size_t size); + + virtual clover::resource &resource(cl_command_queue q); + size_t offset() const; + + clover::root_buffer &parent; + + private: + size_t __offset; + std::map<clover::device *, + std::unique_ptr<clover::sub_resource>> resources; + }; + + struct image : public memory_obj { + protected: + image(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, size_t size, + void *host_ptr); + + public: + virtual clover::resource &resource(cl_command_queue q); + cl_image_format format() const; + size_t width() const; + size_t height() const; + size_t depth() const; + size_t row_pitch() const; + size_t slice_pitch() const; + + private: + cl_image_format __format; + size_t __width; + size_t __height; + size_t __depth; + size_t __row_pitch; + size_t __slice_pitch; + std::map<clover::device *, + std::unique_ptr<clover::root_resource>> resources; + }; + + struct image2d : public image { + public: + image2d(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, size_t width, + size_t height, size_t row_pitch, + void *host_ptr); + + virtual cl_mem_object_type type() const; + }; + + struct image3d : public image { + public: + image3d(clover::context &ctx, cl_mem_flags flags, + const cl_image_format *format, + size_t width, size_t height, size_t depth, + size_t row_pitch, size_t slice_pitch, + void *host_ptr); + + virtual cl_mem_object_type type() const; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/module.cpp b/src/gallium/state_trackers/clover/core/module.cpp new file mode 100644 index 00000000000..1865771443b --- /dev/null +++ b/src/gallium/state_trackers/clover/core/module.cpp @@ -0,0 +1,172 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <type_traits> +#include <algorithm> + +#include "core/module.hpp" + +using namespace clover; + +namespace { + template<typename T, typename = void> + struct __serializer; + + /// Serialize the specified object. + template<typename T> + void + __proc(compat::ostream &os, const T &x) { + __serializer<T>::proc(os, x); + } + + /// Deserialize the specified object. + template<typename T> + void + __proc(compat::istream &is, T &x) { + __serializer<T>::proc(is, x); + } + + template<typename T> + T + __proc(compat::istream &is) { + T x; + __serializer<T>::proc(is, x); + return x; + } + + /// (De)serialize a scalar value. + template<typename T> + struct __serializer<T, typename std::enable_if< + std::is_scalar<T>::value>::type> { + static void + proc(compat::ostream &os, const T &x) { + os.write(reinterpret_cast<const char *>(&x), sizeof(x)); + } + + static void + proc(compat::istream &is, T &x) { + is.read(reinterpret_cast<char *>(&x), sizeof(x)); + } + }; + + /// (De)serialize a vector. + template<typename T> + struct __serializer<compat::vector<T>> { + static void + proc(compat::ostream &os, const compat::vector<T> &v) { + __proc<uint32_t>(os, v.size()); + + for (size_t i = 0; i < v.size(); i++) + __proc<T>(os, v[i]); + } + + static void + proc(compat::istream &is, compat::vector<T> &v) { + v.reserve(__proc<uint32_t>(is)); + + for (size_t i = 0; i < v.size(); i++) + new(&v[i]) T(__proc<T>(is)); + } + }; + + /// (De)serialize a module::section. + template<> + struct __serializer<module::section> { + template<typename S, typename QT> + static void + proc(S &s, QT &x) { + __proc(s, x.type); + __proc(s, x.size); + __proc(s, x.data); + } + }; + + /// (De)serialize a module::argument. + template<> + struct __serializer<module::argument> { + template<typename S, typename QT> + static void + proc(S &s, QT &x) { + __proc(s, x.type); + __proc(s, x.size); + } + }; + + /// (De)serialize a module::symbol. + template<> + struct __serializer<module::symbol> { + template<typename S, typename QT> + static void + proc(S &s, QT &x) { + __proc(s, x.section); + __proc(s, x.offset); + __proc(s, x.args); + } + }; + + /// (De)serialize a module. + template<> + struct __serializer<module> { + template<typename S, typename QT> + static void + proc(S &s, QT &x) { + __proc(s, x.syms); + __proc(s, x.secs); + } + }; +}; + +namespace clover { + void + module::serialize(compat::ostream &os) const { + __proc(os, *this); + } + + module + module::deserialize(compat::istream &is) { + return __proc<module>(is); + } + + const module::symbol & + module::sym(compat::string name) const { + auto it = std::find_if(syms.begin(), syms.end(), [&](const symbol &x) { + return compat::string(x.name) == name; + }); + + if (it == syms.end()) + throw noent_error(); + + return *it; + } + + const module::section & + module::sec(typename section::type type) const { + auto it = std::find_if(secs.begin(), secs.end(), [&](const section &x) { + return x.type == type; + }); + + if (it == secs.end()) + throw noent_error(); + + return *it; + } +} diff --git a/src/gallium/state_trackers/clover/core/module.hpp b/src/gallium/state_trackers/clover/core/module.hpp new file mode 100644 index 00000000000..bc4b203af8e --- /dev/null +++ b/src/gallium/state_trackers/clover/core/module.hpp @@ -0,0 +1,93 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_MODULE_HPP__ +#define __CORE_MODULE_HPP__ + +#include "core/compat.hpp" + +namespace clover { + struct module { + class noent_error { + public: + virtual ~noent_error() {} + }; + + typedef uint32_t resource_id; + typedef uint32_t size_t; + + struct section { + enum type { + text, + data_constant, + data_global, + data_local, + data_private + }; + + resource_id id; + type type; + size_t size; + clover::compat::vector<char> data; + }; + + struct argument { + enum type { + scalar, + constant, + global, + local, + image2d_rd, + image2d_wr, + image3d_rd, + image3d_wr, + sampler + }; + + type type; + size_t size; + }; + + struct symbol { + clover::compat::vector<char> name; + resource_id section; + size_t offset; + clover::compat::vector<argument> args; + }; + + void serialize(compat::ostream &os) const; + static module deserialize(compat::istream &is); + + /// Look up a symbol by name. Throws module::noent_error if not + /// found. + const symbol &sym(compat::string name) const; + + /// Look up a section by type. Throws module::noent_error if not + /// found. + const section &sec(typename section::type type) const; + + clover::compat::vector<symbol> syms; + clover::compat::vector<section> secs; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/program.cpp b/src/gallium/state_trackers/clover/core/program.cpp new file mode 100644 index 00000000000..5ac9f93480e --- /dev/null +++ b/src/gallium/state_trackers/clover/core/program.cpp @@ -0,0 +1,85 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/program.hpp" +#include "core/compiler.hpp" + +using namespace clover; + +_cl_program::_cl_program(clover::context &ctx, + const std::string &source) : + ctx(ctx), __source(source) { +} + +_cl_program::_cl_program(clover::context &ctx, + const std::vector<clover::device *> &devs, + const std::vector<clover::module> &binaries) : + ctx(ctx) { + for_each([&](clover::device *dev, const clover::module &bin) { + __binaries.insert({ dev, bin }); + }, + devs.begin(), devs.end(), binaries.begin()); +} + +void +_cl_program::build(const std::vector<clover::device *> &devs) { + __binaries.clear(); + __logs.clear(); + + for (auto dev : devs) { + try { + auto module = (dev->ir_target() == "tgsi" ? + compile_program_tgsi(__source, dev->ir_target()) : + compile_program_llvm(__source, dev->ir_target())); + __binaries.insert({ dev, module }); + + } catch (build_error &e) { + __logs.insert({ dev, e.what() }); + throw error(CL_BUILD_PROGRAM_FAILURE); + } + } +} + +const std::string & +_cl_program::source() const { + return __source; +} + +const std::map<clover::device *, clover::module> & +_cl_program::binaries() const { + return __binaries; +} + +cl_build_status +_cl_program::build_status(clover::device *dev) const { + return __binaries.count(dev) ? CL_BUILD_SUCCESS : CL_BUILD_NONE; +} + +std::string +_cl_program::build_opts(clover::device *dev) const { + return {}; +} + +std::string +_cl_program::build_log(clover::device *dev) const { + return __logs.count(dev) ? __logs.find(dev)->second : ""; +} diff --git a/src/gallium/state_trackers/clover/core/program.hpp b/src/gallium/state_trackers/clover/core/program.hpp new file mode 100644 index 00000000000..f3858f6ce98 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/program.hpp @@ -0,0 +1,61 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_PROGRAM_HPP__ +#define __CORE_PROGRAM_HPP__ + +#include <map> + +#include "core/base.hpp" +#include "core/context.hpp" +#include "core/module.hpp" + +namespace clover { + typedef struct _cl_program program; +} + +struct _cl_program : public clover::ref_counter { +public: + _cl_program(clover::context &ctx, + const std::string &source); + _cl_program(clover::context &ctx, + const std::vector<clover::device *> &devs, + const std::vector<clover::module> &binaries); + + void build(const std::vector<clover::device *> &devs); + + const std::string &source() const; + const std::map<clover::device *, clover::module> &binaries() const; + + cl_build_status build_status(clover::device *dev) const; + std::string build_opts(clover::device *dev) const; + std::string build_log(clover::device *dev) const; + + clover::context &ctx; + +private: + std::map<clover::device *, clover::module> __binaries; + std::map<clover::device *, std::string> __logs; + std::string __source; +}; + +#endif diff --git a/src/gallium/state_trackers/clover/core/queue.cpp b/src/gallium/state_trackers/clover/core/queue.cpp new file mode 100644 index 00000000000..7e476c715e0 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/queue.cpp @@ -0,0 +1,69 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <algorithm> + +#include "core/queue.hpp" +#include "core/event.hpp" +#include "pipe/p_screen.h" +#include "pipe/p_context.h" + +using namespace clover; + +_cl_command_queue::_cl_command_queue(context &ctx, device &dev, + cl_command_queue_properties props) : + ctx(ctx), dev(dev), __props(props) { + pipe = dev.pipe->context_create(dev.pipe, NULL); + if (!pipe) + throw error(CL_INVALID_DEVICE); +} + +_cl_command_queue::~_cl_command_queue() { + pipe->destroy(pipe); +} + +void +_cl_command_queue::flush() { + pipe_screen *screen = dev.pipe; + pipe_fence_handle *fence = NULL; + + if (!queued_events.empty()) { + // Find out which events have already been signalled. + auto first = queued_events.begin(); + auto last = std::find_if(queued_events.begin(), queued_events.end(), + [](event_ptr &ev) { return !ev->signalled(); }); + + // Flush and fence them. + pipe->flush(pipe, &fence); + std::for_each(first, last, [&](event_ptr &ev) { ev->fence(fence); }); + screen->fence_reference(screen, &fence, NULL); + queued_events.erase(first, last); + } +} + +void +_cl_command_queue::sequence(clover::hard_event *ev) { + if (!queued_events.empty()) + queued_events.back()->chain(ev); + + queued_events.push_back(ev); +} diff --git a/src/gallium/state_trackers/clover/core/queue.hpp b/src/gallium/state_trackers/clover/core/queue.hpp new file mode 100644 index 00000000000..54c949b203f --- /dev/null +++ b/src/gallium/state_trackers/clover/core/queue.hpp @@ -0,0 +1,71 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_QUEUE_HPP__ +#define __CORE_QUEUE_HPP__ + +#include "core/base.hpp" +#include "core/context.hpp" +#include "pipe/p_context.h" + +namespace clover { + typedef struct _cl_command_queue command_queue; + class resource; + class mapping; + class hard_event; +} + +struct _cl_command_queue : public clover::ref_counter { +public: + _cl_command_queue(clover::context &ctx, clover::device &dev, + cl_command_queue_properties props); + _cl_command_queue(const _cl_command_queue &q) = delete; + ~_cl_command_queue(); + + void flush(); + + cl_command_queue_properties props() const { + return __props; + } + + clover::context &ctx; + clover::device &dev; + + friend class clover::resource; + friend class clover::mapping; + friend class clover::hard_event; + friend struct _cl_sampler; + friend struct _cl_kernel; + +private: + /// Serialize a hardware event with respect to the previous ones, + /// and push it to the pending list. + void sequence(clover::hard_event *ev); + + cl_command_queue_properties __props; + pipe_context *pipe; + + typedef clover::ref_ptr<clover::hard_event> event_ptr; + std::vector<event_ptr> queued_events; +}; + +#endif diff --git a/src/gallium/state_trackers/clover/core/resource.cpp b/src/gallium/state_trackers/clover/core/resource.cpp new file mode 100644 index 00000000000..1d241e595aa --- /dev/null +++ b/src/gallium/state_trackers/clover/core/resource.cpp @@ -0,0 +1,192 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/resource.hpp" +#include "pipe/p_screen.h" +#include "util/u_sampler.h" + +using namespace clover; + +namespace { + class box { + public: + box(const resource::point &origin, const resource::point &size) : + pipe({ (unsigned)origin[0], (unsigned)origin[1], + (unsigned)origin[2], (unsigned)size[0], + (unsigned)size[1], (unsigned)size[2] }) { + } + + operator const pipe_box *() { + return &pipe; + } + + protected: + pipe_box pipe; + }; +} + +resource::resource(clover::device &dev, clover::memory_obj &obj) : + dev(dev), obj(obj), pipe(NULL), offset{0} { +} + +resource::~resource() { +} + +void +resource::copy(command_queue &q, const point &origin, const point ®ion, + resource &src_res, const point &src_origin) { + point p = offset + origin; + + q.pipe->resource_copy_region(q.pipe, pipe, 0, p[0], p[1], p[2], + src_res.pipe, 0, + box(src_res.offset + src_origin, region)); +} + +void * +resource::add_map(command_queue &q, cl_map_flags flags, bool blocking, + const point &origin, const point ®ion) { + maps.emplace_back(q, *this, flags, blocking, origin, region); + return maps.back(); +} + +void +resource::del_map(void *p) { + auto it = std::find(maps.begin(), maps.end(), p); + if (it != maps.end()) + maps.erase(it); +} + +unsigned +resource::map_count() const { + return maps.size(); +} + +pipe_sampler_view * +resource::bind_sampler_view(clover::command_queue &q) { + pipe_sampler_view info; + + u_sampler_view_default_template(&info, pipe, pipe->format); + return q.pipe->create_sampler_view(q.pipe, pipe, &info); +} + +void +resource::unbind_sampler_view(clover::command_queue &q, + pipe_sampler_view *st) { + q.pipe->sampler_view_destroy(q.pipe, st); +} + +pipe_surface * +resource::bind_surface(clover::command_queue &q, bool rw) { + pipe_surface info {}; + + info.format = pipe->format; + info.usage = pipe->bind; + info.writable = rw; + + if (pipe->target == PIPE_BUFFER) + info.u.buf.last_element = pipe->width0 - 1; + + return q.pipe->create_surface(q.pipe, pipe, &info); +} + +void +resource::unbind_surface(clover::command_queue &q, pipe_surface *st) { + q.pipe->surface_destroy(q.pipe, st); +} + +root_resource::root_resource(clover::device &dev, clover::memory_obj &obj, + std::string data) : + resource(dev, obj) { + pipe_resource info {}; + + if (image *img = dynamic_cast<image *>(&obj)) { + info.format = translate_format(img->format()); + info.width0 = img->width(); + info.height0 = img->height(); + info.depth0 = img->depth(); + } else { + info.width0 = obj.size(); + } + + info.target = translate_target(obj.type()); + info.bind = (PIPE_BIND_SAMPLER_VIEW | + PIPE_BIND_COMPUTE_RESOURCE | + PIPE_BIND_GLOBAL | + PIPE_BIND_TRANSFER_READ | + PIPE_BIND_TRANSFER_WRITE); + + pipe = dev.pipe->resource_create(dev.pipe, &info); + if (!pipe) + throw error(CL_OUT_OF_RESOURCES); + + assert(data.empty()); // XXX -- initialize it with the supplied data +} + +root_resource::root_resource(clover::device &dev, clover::memory_obj &obj, + clover::root_resource &r) : + resource(dev, obj) { + assert(0); // XXX -- resource shared among dev and r.dev +} + +root_resource::~root_resource() { + dev.pipe->resource_destroy(dev.pipe, pipe); +} + +sub_resource::sub_resource(clover::resource &r, point offset) : + resource(r.dev, r.obj) { + pipe = r.pipe; + offset = r.offset + offset; +} + +mapping::mapping(command_queue &q, resource &r, + cl_map_flags flags, bool blocking, + const resource::point &origin, + const resource::point ®ion) : + pctx(q.pipe) { + unsigned usage = ((flags & CL_MAP_WRITE ? PIPE_TRANSFER_WRITE : 0 ) | + (flags & CL_MAP_READ ? PIPE_TRANSFER_READ : 0 ) | + (blocking ? PIPE_TRANSFER_UNSYNCHRONIZED : 0)); + + pxfer = pctx->get_transfer(pctx, r.pipe, 0, usage, + box(origin + r.offset, region)); + if (!pxfer) + throw error(CL_OUT_OF_RESOURCES); + + p = pctx->transfer_map(pctx, pxfer); + if (!p) { + pctx->transfer_destroy(pctx, pxfer); + throw error(CL_OUT_OF_RESOURCES); + } +} + +mapping::mapping(mapping &&m) : + pctx(m.pctx), pxfer(m.pxfer), p(m.p) { + m.p = NULL; + m.pxfer = NULL; +} + +mapping::~mapping() { + if (pxfer) { + pctx->transfer_unmap(pctx, pxfer); + pctx->transfer_destroy(pctx, pxfer); + } +} diff --git a/src/gallium/state_trackers/clover/core/resource.hpp b/src/gallium/state_trackers/clover/core/resource.hpp new file mode 100644 index 00000000000..d4992972903 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/resource.hpp @@ -0,0 +1,129 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_RESOURCE_HPP__ +#define __CORE_RESOURCE_HPP__ + +#include <list> + +#include "core/base.hpp" +#include "core/memory.hpp" +#include "core/geometry.hpp" +#include "pipe/p_state.h" + +namespace clover { + class mapping; + + /// + /// Class that represents a device-specific instance of some memory + /// object. + /// + class resource { + public: + typedef clover::point<size_t, 3> point; + + resource(const resource &r) = delete; + virtual ~resource(); + + void copy(command_queue &q, const point &origin, const point ®ion, + resource &src_resource, const point &src_origin); + + void *add_map(command_queue &q, cl_map_flags flags, bool blocking, + const point &origin, const point ®ion); + void del_map(void *p); + unsigned map_count() const; + + clover::device &dev; + clover::memory_obj &obj; + + friend class sub_resource; + friend class mapping; + friend struct ::_cl_kernel; + + protected: + resource(clover::device &dev, clover::memory_obj &obj); + + pipe_sampler_view *bind_sampler_view(clover::command_queue &q); + void unbind_sampler_view(clover::command_queue &q, + pipe_sampler_view *st); + + pipe_surface *bind_surface(clover::command_queue &q, bool rw); + void unbind_surface(clover::command_queue &q, pipe_surface *st); + + pipe_resource *pipe; + point offset; + + private: + std::list<mapping> maps; + }; + + /// + /// Resource associated with its own top-level data storage + /// allocated in some device. + /// + class root_resource : public resource { + public: + root_resource(clover::device &dev, clover::memory_obj &obj, + std::string data); + root_resource(clover::device &dev, clover::memory_obj &obj, + root_resource &r); + virtual ~root_resource(); + }; + + /// + /// Resource that reuses a portion of some other resource as data + /// storage. + /// + class sub_resource : public resource { + public: + sub_resource(clover::resource &r, point offset); + }; + + /// + /// Class that represents a mapping of some resource into the CPU + /// memory space. + /// + class mapping { + public: + mapping(command_queue &q, resource &r, cl_map_flags flags, + bool blocking, const resource::point &origin, + const resource::point ®ion); + mapping(const mapping &m) = delete; + mapping(mapping &&m); + ~mapping(); + + operator void *() { + return p; + } + + operator char *() { + return (char *)p; + } + + private: + pipe_context *pctx; + pipe_transfer *pxfer; + void *p; + }; +} + +#endif diff --git a/src/gallium/state_trackers/clover/core/sampler.cpp b/src/gallium/state_trackers/clover/core/sampler.cpp new file mode 100644 index 00000000000..6d683f2b41a --- /dev/null +++ b/src/gallium/state_trackers/clover/core/sampler.cpp @@ -0,0 +1,73 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/sampler.hpp" +#include "pipe/p_state.h" + +using namespace clover; + +_cl_sampler::_cl_sampler(clover::context &ctx, bool norm_mode, + cl_addressing_mode addr_mode, + cl_filter_mode filter_mode) : + ctx(ctx), __norm_mode(norm_mode), + __addr_mode(addr_mode), __filter_mode(filter_mode) { +} + +bool +_cl_sampler::norm_mode() { + return __norm_mode; +} + +cl_addressing_mode +_cl_sampler::addr_mode() { + return __addr_mode; +} + +cl_filter_mode +_cl_sampler::filter_mode() { + return __filter_mode; +} + +void * +_cl_sampler::bind(clover::command_queue &q) { + struct pipe_sampler_state info {}; + + info.normalized_coords = norm_mode(); + + info.wrap_s = info.wrap_t = info.wrap_r = + (addr_mode() == CL_ADDRESS_CLAMP_TO_EDGE ? PIPE_TEX_WRAP_CLAMP_TO_EDGE : + addr_mode() == CL_ADDRESS_CLAMP ? PIPE_TEX_WRAP_CLAMP_TO_BORDER : + addr_mode() == CL_ADDRESS_REPEAT ? PIPE_TEX_WRAP_REPEAT : + addr_mode() == CL_ADDRESS_MIRRORED_REPEAT ? PIPE_TEX_WRAP_MIRROR_REPEAT : + PIPE_TEX_WRAP_CLAMP_TO_EDGE); + + info.min_img_filter = info.mag_img_filter = + (filter_mode() == CL_FILTER_LINEAR ? PIPE_TEX_FILTER_LINEAR : + PIPE_TEX_FILTER_NEAREST); + + return q.pipe->create_sampler_state(q.pipe, &info); +} + +void +_cl_sampler::unbind(clover::command_queue &q, void *st) { + q.pipe->delete_sampler_state(q.pipe, st); +} diff --git a/src/gallium/state_trackers/clover/core/sampler.hpp b/src/gallium/state_trackers/clover/core/sampler.hpp new file mode 100644 index 00000000000..5bb5bccb1a1 --- /dev/null +++ b/src/gallium/state_trackers/clover/core/sampler.hpp @@ -0,0 +1,55 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#ifndef __CORE_SAMPLER_HPP__ +#define __CORE_SAMPLER_HPP__ + +#include "core/base.hpp" +#include "core/queue.hpp" + +namespace clover { + typedef struct _cl_sampler sampler; +} + +struct _cl_sampler : public clover::ref_counter { +public: + _cl_sampler(clover::context &ctx, bool norm_mode, + cl_addressing_mode addr_mode, cl_filter_mode filter_mode); + + bool norm_mode(); + cl_addressing_mode addr_mode(); + cl_filter_mode filter_mode(); + + clover::context &ctx; + + friend class _cl_kernel; + +private: + void *bind(clover::command_queue &q); + void unbind(clover::command_queue &q, void *st); + + bool __norm_mode; + cl_addressing_mode __addr_mode; + cl_filter_mode __filter_mode; +}; + +#endif diff --git a/src/gallium/state_trackers/clover/llvm/invocation.cpp b/src/gallium/state_trackers/clover/llvm/invocation.cpp new file mode 100644 index 00000000000..89e21bf9289 --- /dev/null +++ b/src/gallium/state_trackers/clover/llvm/invocation.cpp @@ -0,0 +1,94 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include "core/compiler.hpp" + +#if 0 +#include <clang/Frontend/CompilerInstance.h> +#include <clang/Frontend/TextDiagnosticPrinter.h> +#include <clang/CodeGen/CodeGenAction.h> +#include <llvm/LLVMContext.h> +#include <llvm/Support/TargetSelect.h> +#include <llvm/Support/MemoryBuffer.h> + +#include <iostream> +#include <iomanip> +#include <fstream> +#include <cstdio> +#endif + +using namespace clover; + +#if 0 +namespace { + void + build_binary(const std::string &source, const std::string &target, + const std::string &name) { + clang::CompilerInstance c; + clang::EmitObjAction act(&llvm::getGlobalContext()); + std::string log; + llvm::raw_string_ostream s_log(log); + + LLVMInitializeTGSITarget(); + LLVMInitializeTGSITargetInfo(); + LLVMInitializeTGSITargetMC(); + LLVMInitializeTGSIAsmPrinter(); + + c.getFrontendOpts().Inputs.push_back( + std::make_pair(clang::IK_OpenCL, name)); + c.getHeaderSearchOpts().UseBuiltinIncludes = false; + c.getHeaderSearchOpts().UseStandardIncludes = false; + c.getLangOpts().NoBuiltin = true; + c.getTargetOpts().Triple = target; + c.getInvocation().setLangDefaults(clang::IK_OpenCL); + c.createDiagnostics(0, NULL, new clang::TextDiagnosticPrinter( + s_log, c.getDiagnosticOpts())); + + c.getPreprocessorOpts().addRemappedFile( + name, llvm::MemoryBuffer::getMemBuffer(source)); + + if (!c.ExecuteAction(act)) + throw build_error(log); + } + + module + load_binary(const char *name) { + std::ifstream fs((name)); + std::vector<unsigned char> str((std::istreambuf_iterator<char>(fs)), + (std::istreambuf_iterator<char>())); + compat::istream cs(str); + return module::deserialize(cs); + } +} +#endif + +module +clover::compile_program_llvm(const compat::string &source, + const compat::string &target) { +#if 0 + build_binary(source, target, "cl_input"); + module m = load_binary("cl_input.o"); + std::remove("cl_input.o"); + return m; +#endif + return module(); +} diff --git a/src/gallium/state_trackers/clover/tgsi/compiler.cpp b/src/gallium/state_trackers/clover/tgsi/compiler.cpp new file mode 100644 index 00000000000..eb27db1aa76 --- /dev/null +++ b/src/gallium/state_trackers/clover/tgsi/compiler.cpp @@ -0,0 +1,100 @@ +// +// Copyright 2012 Francisco Jerez +// +// Permission is hereby granted, free of charge, to any person obtaining a +// copy of this software and associated documentation files (the "Software"), +// to deal in 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: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// 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 AUTHORS 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 IN THE +// SOFTWARE. +// + +#include <sstream> + +#include "core/compiler.hpp" + +#include "tgsi/tgsi_parse.h" +#include "tgsi/tgsi_text.h" +#include "util/u_memory.h" + +using namespace clover; + +namespace { + void + read_header(const std::string &header, module &m) { + std::istringstream ls(header); + std::string line; + + while (getline(ls, line)) { + std::istringstream ts(line); + std::string name, tok; + module::size_t offset; + compat::vector<module::argument> args; + + if (!(ts >> name)) + continue; + + if (!(ts >> offset)) + throw build_error("invalid kernel start address"); + + while (ts >> tok) { + if (tok == "scalar") + args.push_back({ module::argument::scalar, 4 }); + else if (tok == "global") + args.push_back({ module::argument::global, 4 }); + else if (tok == "local") + args.push_back({ module::argument::local, 4 }); + else if (tok == "constant") + args.push_back({ module::argument::constant, 4 }); + else if (tok == "image2d_rd") + args.push_back({ module::argument::image2d_rd, 4 }); + else if (tok == "image2d_wr") + args.push_back({ module::argument::image2d_wr, 4 }); + else if (tok == "image3d_rd") + args.push_back({ module::argument::image3d_rd, 4 }); + else if (tok == "image3d_wr") + args.push_back({ module::argument::image3d_wr, 4 }); + else if (tok == "sampler") + args.push_back({ module::argument::sampler, 0 }); + else + throw build_error("invalid kernel argument"); + } + + m.syms.push_back({ name, 0, offset, args }); + } + } + + void + read_body(const char *source, module &m) { + tgsi_token prog[1024]; + + if (!tgsi_text_translate(source, prog, Elements(prog))) + throw build_error("translate failed"); + + unsigned sz = tgsi_num_tokens(prog) * sizeof(tgsi_token); + m.secs.push_back({ 0, module::section::text, sz, { (char *)prog, sz } }); + } +} + +module +clover::compile_program_tgsi(const compat::string &source, + const compat::string &target) { + const char *body = source.find("COMP\n"); + module m; + + read_header({ source.begin(), body }, m); + read_body(body, m); + + return m; +} diff --git a/src/gallium/targets/opencl/Makefile.am b/src/gallium/targets/opencl/Makefile.am new file mode 100644 index 00000000000..0d233c11b8d --- /dev/null +++ b/src/gallium/targets/opencl/Makefile.am @@ -0,0 +1,36 @@ +AUTOMAKE_OPTIONS = subdir-objects + +lib_LTLIBRARIES = libOpenCL.la + +libOpenCL_la_LDFLAGS = \ + -version-number 1:0 + +libOpenCL_la_LIBADD = \ + $(top_builddir)/src/gallium/state_trackers/clover/libclover.la \ + $(top_builddir)/src/gallium/auxiliary/libgallium.a \ + $(GALLIUM_PIPE_LOADER_LIBS) $(LIBUDEV_LIBS) \ + -ldl + +libOpenCL_la_SOURCES = + +# Force usage of a C++ linker +nodist_EXTRA_libOpenCL_la_SOURCES = dummy.cpp + +PIPE_SRC_DIR = $(top_srcdir)/src/gallium/targets/pipe-loader + +# Provide compatibility with scripts for the old Mesa build system for +# a while by putting a link to the driver into /lib of the build tree. +all-local: libOpenCL.la + @$(MAKE) -C $(PIPE_SRC_DIR) + $(MKDIR_P) $(top_builddir)/$(LIB_DIR) + ln -f .libs/libOpenCL.so* $(top_builddir)/$(LIB_DIR)/ + +install-exec-local: + @$(MAKE) -C $(PIPE_SRC_DIR) PIPE_INSTALL_DIR=$(OPENCL_LIB_INSTALL_DIR) install + +clean-local: + @$(MAKE) -C $(PIPE_SRC_DIR) clean + +# FIXME: Remove when the rest of Gallium is converted to automake. +TOP=$(top_builddir) +default: all |