KhronosGroup / OpenCL-CTS

The OpenCL Conformance Tests
Apache License 2.0
185 stars 197 forks source link

Dump compilation options list within a kernel file as a comment #1312

Open azabazno opened 3 years ago

azabazno commented 3 years ago

The idea is to have something like:

// OPTIONS: -cl-std=CL3.0
__kernel void test0(void) { printf("%-+5d\n", 10); }

when, for example, option --compilation-cache-mode dump-cl-files or other special option is used. This is needed for maintaining OpenCL-CTS-Kernel-Dumps repo (see discussion here).

bashbaug commented 3 years ago

An alternative to consider is the method used by the OpenCL Intercept Layer, which emits the program source in one file and the program options in a different file. See, for example the description of DumpProgramSource:

If set to a nonzero value, the Intercept Layer for OpenCL Applications will dump every string passed to clCreateProgramWithSource() to its own file. The filename will have the form "CLI_\<Program Number>_\<Unique Program Hash Code>_source.cl". Program options will be dumped to the same directory with the filename "CLI_\<Program Number>_\<Unique Program Hash Code>_\<Compile Count>_\<Unique Build Options Hash Code>_\<API>_options.txt", where API is an empty string for clBuildProgram(), "compile" for clCompileProgram(), and "link" for clLinkProgram().

While I believe both options could work, dumping to a separate file has several advantages:

  1. No need to do string manipulation to extract the options from the single kernel dump. (To be fair, some amount of string manipulation is required to match the options file to its source file.)
  2. More important: Separating into different files works for other methods of creating and compiling program objects, such as program objects created from source, or program objects that are separately compiled and linked.
kosarev commented 3 years ago

Note that using the intercepting layer would make the source-dumping facility to be unnecessarily dependent on it whereas it seems for the task of dumping test kernels calling the code generators and some file functions should be enough. (https://github.com/KhronosGroup/OpenCL-CTS-Kernel-Dumps/pull/1 discusses this in more detail.)

Re string manipulations: that implies loading another (potentially missing) file is easier than skipping the // OPTIONS: in the beginning of the first line? For multi-object inputs, isn't it the case that to support them we will need to know the file names and probably some command-line flags, which again sounds like // OPTIONS:? So the main test file with the flags plus whatever supplementary sources and binaries needed.

I think it would be fair to mention that keeping flags in the source makes the most of the test files self-consistent. May sound important to some. That's aside from things like not doubling the number of files to deal with.

bashbaug commented 3 years ago

using the intercepting layer would make the source-dumping facility to be unnecessarily dependent on it

To be clear, I am NOT suggesting to use the OpenCL Intercept Layer as the source dumping facility, especially for CTS kernels.

I only mention it because it has a solution to this problem that is implemented and widely used, and to describe some of the rationale that guided us towards this solution.

kosarev commented 3 years ago

@bashbaug Hi Ben, https://github.com/KhronosGroup/OpenCL-CTS-Kernel-Dumps/pull/1 is about to be landed. Are we fine with keeping // OPTIONS: within the sources?