| Name Strings |
| |
| cl_altera_compiler_mode |
| |
| Contributors |
| |
| David Neto, Altera Corporation |
| |
| Contact |
| |
| Michael Kinsner, mkinsner 'at' altera 'dot' com |
| |
| IP Status |
| |
| No known IP claims. |
| |
| Version |
| |
| Version 1, 2014-02-06 |
| |
| Number |
| |
| OpenCL Extension #30 |
| |
| Status |
| |
| Complete. |
| Shipping with the Altera SDK for OpenCL, version 14.0 |
| |
| An earlier version shipped with the Altera SDK for OpenCL, version 13.0 |
| |
| Extension Type |
| |
| OpenCL platform extension |
| |
| Dependencies |
| |
| OpenCL 1.0 is required. |
| This document is written against revision 48 of the OpenCL 1.0 specification. |
| |
| Overview |
| |
| This extension specifies alternative device code compilation flows |
| for OpenCL. |
| |
| In standard OpenCL, there are two methods for preparing code for |
| execution on an OpenCL device. The first is to compile OpenCL C code |
| from source text using an online compiler. The second is to load a |
| precompiled device binary. |
| |
| This extension specifies mechanisms to enable alternative workflows for |
| compiling and loading device code. They enable simplified or optimized |
| support for the following development and deployment scenarios: |
| |
| 1. The target OpenCL device is not available in the development |
| environment. |
| |
| 2. The platform does not provide an online compiler for the target |
| OpenCL device, or full online compilation is not desirable during |
| host application development. |
| |
| 3. The application only requires one device program, and does not |
| need read access to the device binary. |
| |
| The development and deployment of a particular application may involve |
| more than one of these scenarios. |
| |
| The mechanisms specified by this extension are: |
| |
| 1. An "offline device" feature: This is the ability to specify that |
| the OpenCL runtime should partially emulate the presence a device. |
| From the perspective of the host program, emulation should be |
| complete, except that kernel code may have no effect. For example, |
| kernels are enqueued, buffers are copied or mapped as required, |
| and event profiling information is updated to reflect command |
| progression. However side effects produced by running the code |
| inside kernels may not be visible. |
| |
| 2. A "create program executable library" compiler mode: |
| The clBuildProgram API method does not fully compile device code. |
| Instead, it performs an online stub compilation, and saves enough |
| information in a data store to perform a deferred full (offline) |
| compilation. |
| |
| 2a. The stub compilation produces a binary which includes enough |
| kernel interface information to satisfy further host activities |
| such as creating cl_kernel objects, setting kernel arguments, |
| and enqueueing kernels for execution. But the stub binary |
| may not have executable code, so the execution of kernels from |
| such a binary may not produce valid data. |
| |
| 2b. The data store is used to perform offline compilation of |
| kernel programs. Each logically distinct call of clBuildProgram |
| generates an entry in the data store. Each entry is a tuple |
| consisting of: |
| |
| - The target device |
| |
| - The kernel program source |
| |
| - The options supplied to clBuildProgram |
| |
| - A set of instructions (such as a script) for offline |
| compilation of the kernel program. The result of executing |
| these instructions is to associate a device binary (including |
| executable device code) with this tuple in the data store. |
| |
| 3. A "use program executable library" mode: We assume that an |
| offline compilation has been performed for all entries in the |
| data store generated by use of the "create program executable |
| library" mode. In the "use program exectuable library" mode, |
| the clBuildProgram API method performs a lookup in the data store |
| based on the following attributes: |
| |
| - The target device |
| |
| - The kernel program source |
| |
| - The options supplied to clBuildProgram |
| |
| If such an entry exists in the data store, then the associated |
| binary is loaded into the cl_program object. If no such entry |
| exists, or if the entry is not associated with a device binary, |
| then an error is returned. |
| |
| 4. A "preloaded binary only" compiler mode: In this mode, the |
| runtime does not compile any programs, does not load the device |
| with new code at runtime, and does not necessarily produce a valid |
| device binary when requested by the CL_PROGRAM_BINARIES query to |
| the clGetProgramInfo API method. Instead: |
| |
| - The device is assumed to already have a program loaded |
| before the cl_context is created. |
| |
| - All cl_program objects behave as if they refer to the preloaded |
| binary |
| |
| - The clBuildProgram API method is a no-op except for |
| setting a successful build status. |
| |
| The following paragraphs describe the benefits provided by the |
| mechanisms specified in this extension. |
| |
| Scenario 1: Developing the host part of an OpenCL application without |
| a device being present: |
| |
| The offline device feature enables the execution of a host program |
| even while a device is not present in the system. This is |
| useful for developing an application before a device has been |
| manufactured, is otherwise scarce, or unavailable. Although the |
| data coming back from the emulated device is invalid, host code |
| interaction with the runtime may be developed and tested. |
| |
| Scenario 2: Porting an existing application to a platform without an |
| online compiler: |
| |
| Many OpenCL platforms provide an online OpenCL C compiler |
| for each device. Therefore host applications often use |
| clCreateProgramWithSource and clBuildProgram to create |
| executable device code. The standard method of porting such |
| a program to an environment without an online compiler is to |
| replace the compile-from-source sequence of API calls with a |
| find-and-load-from-binary code sequence. Transforming the host |
| application in this way introduces platform dependencies (such |
| as APIs to access a filesystem) and adds bookkeeping complexity |
| to map the original OpenCL C source and compile options with |
| an associated device binary. The porting and verification effort |
| can be quite onerous if many portions of the host code must be |
| modified, or if the device source code is parameterized at runtime. |
| |
| We can reduce porting effort in this scenario by using a three-phase |
| approach. |
| |
| In the first phase, we run the host program with the runtime |
| configured to use offline device mode, and also the "create program |
| executable libary" compiler mode. This generates a data store |
| (the program executable library) containing enough information |
| to compile the required device programs in an offline manner, |
| i.e. outside the control of the host program. |
| |
| In the second phase, we perform an offline compilation of all |
| entries in the data store. This is done by enumerating the entries |
| in the data store and following the compilation instructions for |
| each entry. |
| |
| In the third phase, the application is fully functional. The host |
| is configured to use "use program executable library" compiler mode. |
| Any API request to compile a device program from source is translated |
| into a lookup of the (fully functional) device binary in the data store. |
| |
| This approach works when: |
| |
| - The device has a stable name. That is, the device name does not |
| change from one run to another. |
| |
| - The host program calls clBuildProgram with only a finite |
| set of combinations of device, kernel source, and build |
| options. |
| |
| - Those combinations are consistent from one run to the next. |
| That is, the kernel source and build options do not depend |
| on the data results from executing a previous kernel. |
| |
| - Other context upon which the compilation depends is stable. |
| For example, the location and contents of files included |
| from the OpenCL C source remains the same. |
| |
| Scenario 3: An application needs only one cl_program, runs in a |
| constrained environment, and may need "instant on" behaviour: |
| |
| The "preloaded binary only" compiler mode is desiged to be used |
| by embedded applications. Such applications typically have |
| the following constraints: |
| |
| - The application is simple enough that only one cl_program |
| is required. |
| |
| - The target system has tight memory and time constraints. |
| It may be infeasible or undesirable to process a full device |
| binary through the standard OpenCL runtime APIs. For example, |
| the device binary might be too large to map into host memory, |
| or the application cannot tolerate the delay to load the |
| device binary from external storage. |
| |
| - The application has no requirement to recover the device |
| binary via the CL_PROGRAM_BINARIES query to clGetProgramInfo. |
| Even moreso, it may be undesirable for intellectual property |
| reasons or otherwise to allow the host program to recover |
| the device binary. |
| |
| - It is desirable to develop the application in the normal way, |
| but to minimize host program code changes before deployment. |
| That is, application development uses the standard |
| clCreateProgramWithSource, or clCreateProgramWithBinary APIs. |
| But the transition to a test or production environment |
| is limited to a configuration change, as opposed to a code |
| change. |
| |
| Header File |
| |
| Interface constants are defined in cl_ext.h |
| |
| New Tokens |
| |
| New context properties to be used in the <properties> array |
| supplied to to clCreateContext or clCreateContextFromType: |
| |
| CL_CONTEXT_COMPILER_MODE_ALTERA 0x40F0 |
| CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_ALTERA 0x40F1 |
| CL_CONTEXT_OFFLINE_DEVICE_ALTERA 0x40F2 |
| |
| Values to be supplied for context property |
| CL_CONTEXT_COMPILER_MODE_ALTERA: |
| |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_ALTERA 0 |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA 1 |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_USE_EXE_LIBRARY_ALTERA 2 |
| CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_ALTERA 3 |
| |
| |
| |
| Additions to Chapter 4 of the OpenCL 1.0 (v48) Specification |
| |
| Add the following to Table 4.4, List of supported properties by |
| clCreateContext: |
| |
| |
| cl_context_properties enum: |
| CL_CONTEXT_OFFLINE_DEVICE_ALTERA |
| Property value: |
| const char* |
| Description: |
| Specifies that the runtime should partly emulate the presence |
| of the named device. The device should behave normally except |
| that kernels executed on such an "offline" device might |
| not produce any side effects. |
| |
| The specified name value should be the first word in the full |
| name returned by the CL_DEVICE_NAME, in C-style form. |
| |
| An implementation may restrict the use of this property as |
| follows: |
| |
| - A platform may require external initialization when using |
| this context property. For example, Altera's |
| platform implementations require that environment variable |
| CL_CONTEXT_OFFLINE_DEVICE_ALTERA be set to the same device |
| name string as supplied to this property. |
| |
| - When this property is specified, the specified device |
| may be the only available device in the platform. |
| |
| This property may be useful for developing or porting |
| applications when no online compiler is available, and |
| when the CL_CONTEXT_COMPILER_MODE_ALTERA property is set to |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA. |
| |
| |
| cl_context_properties enum: |
| CL_CONTEXT_COMPILER_MODE_ALTERA |
| Property value: |
| cl_ulong |
| Description: |
| For devices without an online compiler, this property |
| specifies alternative behaviour for building, loading, |
| and querying device programs. |
| (Platforms implementing the embedded profile are the |
| only ones that may omit providing an online compiler.) |
| |
| The value should be one of the following: |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_ALTERA, |
| CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_ALTERA, |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA, |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_USE_EXE_LIBRARY_ALTERA. |
| |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_ALTERA - The clBuildProgram |
| and clCreateProgramWithBinary, and clGetProgramInfo |
| API methods behave in the standard way. |
| In particular, clBuildProgram may fail with error |
| CL_COMPILER_NOT_AVAILABLE. |
| This value is the default, and is the only value |
| which specifies conformant behaviour. |
| |
| CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_ALTERA - The device |
| is assumed to already have a program loaded before |
| the first OpenCL APIs are invoked. |
| All cl_program objects behave as if they refer to the |
| preloaded device binary. |
| The clBuildProgram API method always succeeds when compiling |
| for this device, but does not produce an executable |
| device binary. |
| The clCreateProgramWithBinary does not validate |
| the provided binary. It always succeeds provided its |
| arguments are well-formed. |
| The data returned by the CL_PROGRAM_BINARIES query |
| for clGetProgramInfo is unspecified. It may not be |
| a valid device binary. |
| |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA - |
| When compiling a program created with |
| clCreateProgramWithSource, the clBuildProgram method |
| performs a "stub" compilation, and populates a data store |
| with enough information to perform a program compilation |
| offline (i.e. outside the control of the host program). |
| |
| The online stub compilation produces a binary which includes |
| enough information to satsify further runtime API activities |
| such as creating cl_kernel objects, setting kernel arguments, |
| and enqueing kernels for execution. The stub binary may not |
| have executable code, so the execution of kernels from such |
| a binary may not produce valid data. |
| |
| This value for CL_CONTEXT_COMPILER_MODE_ALTERA is designed |
| to be used with the CL_CONTEXT_OFFLINE_DEVICE_ALTERA context |
| property. |
| |
| Additionally, each logically distinct call of clBuildProgram |
| generates an entry in an external data store which persists |
| beyond the lifetime of the host program. Each entry is a |
| tuple consisting of: |
| |
| - The target device |
| |
| - The kernel program source |
| |
| - The options supplied to clBuildProgram |
| |
| - A set of instructions (such as a script) for offline |
| compilation of the kernel program. The result of executing |
| these instructions is to associate a device binary (including |
| executable device code) with this tuple in the data store. |
| |
| In Altera's platform implementation, the data store is a tree of |
| directories and files. Each data store entry is a leaf |
| directory in this tree containing: |
| - a file named "kernels.cl" containing the kernel program |
| source |
| - a file named "build.cmd" containing the list of operating |
| system commands to be used to compile the program to a |
| device binary, and to store that device binary in a file |
| named "kernels.aocx" in the same directory. |
| |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_USE_EXE_LIBRARY_ALTERA - |
| When compiling a program created with clCreateProgramWithSoruce, |
| the clBuildProgram API method performs a lookup in the data store |
| based on the following attributes: |
| |
| - The target device |
| |
| - The kernel program source |
| |
| - The options supplied to clBuildProgram |
| |
| If such an entry exists in the data store, then the associated |
| binary is loaded into the cl_program. If no such entry exists, |
| or if the entry is not associated with a device binary, then an |
| error is returned. |
| |
| |
| cl_context_properties enum: |
| CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_ALTERA |
| Property value: |
| const char* |
| Description: |
| Specifies the filesystem root directory |
| for the data store used when either value |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA or value |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_USE_EXE_LIBRARY_ALTERA is specified |
| for context property CL_CONTEXT_COMPILER_MODE_ALTERA. |
| |
| The specified value may be a relative directory name, but |
| will be resolved to an absolute path at context creation time. |
| |
| If this property is left unspecified, then data store root |
| is implementation-defined. |
| |
| In Altera's platform implementation, the default is the |
| "aocl_program_library" subdirectory in the current directory |
| in effect at context creation time. |
| |
| |
| |
| Additions to Chapter 5 of the OpenCL 1.0 (v48) Specification |
| |
| Additions to Section 5.4.1 Creating Program Objects |
| |
| When context property CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA is |
| specified with value CL_CONTEXT_COMPILER_PRELOADED_BINARY_ONLY_ALTERA, |
| the clCreateProgramWithBinary API method behaves in a non-standard |
| way. See the description of this property value in Table 4.4 for |
| more information. |
| |
| Additions to Section 5.4.2 Building Program Exectuables |
| |
| When context property CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA is |
| specified, the clBuildProgram API method behaves in a non-standard |
| way when compiling for devices without an online compiler. |
| See the description of this property value in Table 4.4 for more |
| information. |
| |
| Additions to Section 5.4.5 Program Object Queries |
| |
| The following is added to the description of the |
| CL_PROGRAM_BINARIES query in Table 5.11: |
| |
| When context property CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA |
| is specified with either value |
| CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_ALTERA |
| or value CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA, |
| then the binary returned by this query may not be a valid device |
| binary. |
| See the description of the CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA |
| property value in Table 4.4 for more information. |
| |
| Additions to Section 5.6 Executing Kernels |
| |
| When context property CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA |
| is specified with value |
| CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA, |
| or if context property CL_CONTEXT_OFFLINE_DEVICE_ALTERA is specified, |
| then a kernel execution may behave as if the kernel code |
| produces no side effects, i.e. as if the body of the kernel |
| contains no statements. |
| (All other runtime activites must occur. For example, |
| the kernels are still enqueued, buffers are copied or mapped |
| as required to satisfy kernel arguments, and event profiling |
| information is updated to reflect command progression.) |
| See the description of the CL_CONTEXT_OFFLINE_COMPILER_MODE_ALTERA |
| property value in Table 4.4 for more information. |
| |
| |
| Implementation Notes |
| |
| For each context property defined in this extension specification, |
| Altera's platform implementation allows the property to be specified |
| by setting an environment variable of the same name. |
| |
| For CL_CONTEXT_OFFLINE_DEVICE_ALTERA, the environment variable |
| of *must* be set before any platform APIs are called. |
| |
| For CL_CONTEXT_COMPILER_MODE_ALTERA, value of the environment variable |
| is the numerical value for the corresponding value enum as defined |
| in the CL/cl_ext.h header file. |
| |
| Precedence: For context properties CL_CONTEXT_COMPILER_MODE_ALTERA and |
| CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_ALTERA, any value explicitly |
| provided to the clCreateContext or clCreateContextFromType |
| API methods will override an environment variable setting. |
| |
| If property CL_CONTEXT_OFFLINE_DEVICE_ALTERA is explicitly |
| provided to the clCreateCOTNEXT or clCreateContextFromType |
| methods, then it must match the value set for environment variable |
| CL_CONTEXT_OFFLINE_DEVICE_ALTERA. |
| |
| |
| Issues |
| |
| 1. The use of context property CL_CONTEXT_OFFLINE_DEVICE_ALTERA is |
| rather restricted: |
| |
| - It must be set at the "beginning of time" for a host application, |
| e.g. via an environment variable setting. |
| |
| - If an "offline device" is used in a context, then online |
| devices cannot be used in any context for the platform. |
| |
| In this light, it might be more natural to call it a "platform" |
| property. But in OpenCL platforms are stateless, so this would |
| not be appropriate. Instead the most basic runtime controls are |
| applied to contexts, not platforms. |
| |
| These restrictions reflect limitations of Altera's original |
| implementation. However, the feature in its current form is still |
| quite useful for application development and porting. |
| |
| |
| Sample Code |
| |
| Example for using an "offline device", with creation and initialization of |
| a data store for offline compilation of program binaries. |
| |
| #include <CL/opencl.h> |
| #include <stdlib.h> |
| #include <assert.h> |
| #define CHECK(X) assert(CL_SUCCESS == (X)) |
| |
| int main(...) { |
| cl_platform platform = 0; |
| cl_device_id device = 0; |
| cl_context context = 0; |
| cl_int status = 0; |
| |
| // Specify an offline device via environment variable here, or externally |
| // before program startup. |
| // This must occur befor the first OpenCL API method call. |
| setenv("CL_CONTEXT_OFFLINE_DEVICE_ALTERA","mydevice"); |
| |
| CHECK( clGetPlatformIDs(1,&platform,0) ); |
| CHECK( clGetDeviceIDs(platform,CL_DEVICE_TYPE_ACCELERATOR,1,&device,0) ); |
| |
| cl_context_properties props[] = { |
| |
| // Specify creation and initialization of a program library data store. |
| CL_CONTEXT_COMPILER_MODE_ALTERA, |
| (cl_context_properties)CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_ALTERA |
| |
| // Where is the data store rooted? |
| CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_ALTERA |
| (cl_context_properties)"/data/myproject/proglib" |
| |
| // Terminate the properties list. |
| 0 |
| }; |
| context = clCreateContext( props, 1, &device, 0, 0, &status ); |
| CHECK( status ); |
| |
| cl_command_queue cq = clCreateCommandQueue( context, device, 0, &status ); |
| CHECK( status ); |
| |
| const char* source = "kernel void foo( global int* A ) { *A = 42; }"; |
| cl_program program = clCreateProgramWithSource( context, 1, &source, NULL, &status ); |
| CHECK( status ); |
| |
| // Perform a stub compilation, and create an entry for this (device,source,options) |
| // combination in the program library data store. |
| CHECK( clBuildProgram( program, 1, &device, "-cl-opt-disable", 0, 0 ) ); |
| |
| cl_mem mem = clCreateBuffer(context,CL_MEM_READ_WRITE,sizeof(cl_int),0,&status); |
| CHECK( status ); |
| |
| // Only a stub compilation has been performed. |
| // But this still enables full setup and enqueue of a kernel for execution. |
| cl_kernel kernel = clCreateKernel( program, "foo", &status ); |
| CHECK( status ); |
| CHECK( clSetKernelArg( kernel, 0, sizeof(cl_mem), &mem ) ); |
| |
| // Expect that kernels scheduled by clEnqueueNDRangeKernel and |
| // clEnqueueTask may not produce expected side effects. |
| // It may appear that each kernel executes no statements. |
| |
| CHECK( clEnqueueTask( cq, kernel, 0, 0, 0 ) ); |
| cl_int the_answer = 0; |
| CHECK( clEnqueueReadBuffer( cq, mem, 1 /*block*/, 0, sizeof(the_answer),&the_answer,0,0,0) ); |
| |
| // Because we're using an "offline device", we can't rely on the_answer being 42. |
| |
| // ... |
| |
| |
| |
| Example for using an already-populated data store of program binaries compiled |
| in an offline manner. The code is the same as the previous example, except |
| use these context properties instead: |
| |
| cl_context_properties props[] = { |
| |
| // Specify the use of a program data store library of pre-compiled device binaries. |
| CL_CONTEXT_COMPILER_MODE_ALTERA, |
| (cl_context_properties)CL_CONTEXT_COMPILER_MODE_OFFLINE_USE_EXE_LIBRARY_ALTERA |
| |
| // Where is the data store rooted? |
| CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_ALTERA |
| (cl_context_properties)"/data/myproject/proglib" |
| |
| // Terminate the properties list. |
| 0 |
| }; |
| |
| |
| |
| Example for using a "preloaded binary": |
| |
| cl_context_properties props[] = { |
| |
| // Specify the use of a preloaded binary. |
| CL_CONTEXT_COMPILER_MODE_ALTERA, |
| (cl_context_properties)CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_ALTERA, |
| |
| // Terminate the properties list. |
| 0 |
| }; |
| context = clCreateContext( props, 1, &device, 0, 0, &status ); |
| |
| ... |
| |
| // All cl_program objects refer to the preloaded binary. |
| // clCreateProgramWithBinary accepts invalid device binaries. |
| // clBuildProgram succeeds but does not actually compile the program from source. |
| |
| |
| Conformance Tests |
| |
| None. |
| |
| Revision History |
| |
| Version 0, 2013-05-06 - Initial revision. |
| Documents the behaviours in the Altera SDK for OpenCL, version 13.0. |
| |
| Version 1, 2014-02-06 - Assigned registered values to tokens. |
| |
