212 lines
8.7 KiB
Markdown
212 lines
8.7 KiB
Markdown
# CLProgramCreate
|
||
|
||
Creates an OpenCL program from a source code.
|
||
|
||
```
|
||
int CLProgramCreate(
|
||
int context, // Handle to an OpenCL context
|
||
const string source // Source code
|
||
);
|
||
|
||
```
|
||
|
||
An overloaded function version creates an OpenCL program and writes compiler messages into the passed string.
|
||
|
||
```
|
||
int CLProgramCreate(
|
||
int context, // Handle to an OpenCL context
|
||
const string source, // Source code
|
||
string &build_log // A string for receiving the compilation log
|
||
);
|
||
|
||
```
|
||
|
||
Parameters
|
||
|
||
context
|
||
|
||
[in] Handle of the OpenCL context.
|
||
|
||
source
|
||
|
||
[in] String with the source code of the OpenCL program.
|
||
|
||
&build_log
|
||
|
||
[in] A string for receiving the OpenCL compiler messages.
|
||
|
||
Return Value
|
||
|
||
A handle to an OpenCL object if successful. In case of error -1 is returned. For information about the error, use the [GetLastError()](/en/docs/check/getlasterror) function.
|
||
|
||
Note
|
||
|
||
At the moment, the following error codes are used:
|
||
|
||
- ERR_OPENCL_INVALID_HANDLE – invalid handle to the context OpenCL.
|
||
- ERR_INVALID_PARAMETER – invalid string parameter.
|
||
- ERR_NOT_ENOUGH_MEMORY – not enough memory to complete operation.
|
||
- ERR_OPENCL_PROGRAM_CREATE – internal error of OpenCL or compilation error.
|
||
|
||
In some graphic cards working with the [double](/en/docs/basis/types/double) type numbers is disabled by default. This can lead to compilation error 5105. To enable support for the double type numbers, please add the following directive to your OpenCL program: [#pragma OPENCL EXTENSION cl_khr_fp64 : enable](https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/cl_khr_fp64.html). However if a graphic card doesn't support double, enabling this directive won't be of help.
|
||
|
||
Example:
|
||
|
||
```
|
||
//+------------------------------------------------------------------+
|
||
//| OpenCL kernel |
|
||
//+------------------------------------------------------------------+
|
||
const string
|
||
cl_src=
|
||
//--- by default some GPU doesn't support doubles
|
||
//--- cl_khr_fp64 directive is used to enable work with doubles
|
||
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable \r\n"
|
||
//--- OpenCL kernel function
|
||
"__kernel void Test_GPU(__global double *data, \r\n"
|
||
" const int N, \r\n"
|
||
" const int total_arrays) \r\n"
|
||
" { \r\n"
|
||
" uint kernel_index=get_global_id(0); \r\n"
|
||
" if (kernel_index>total_arrays) return; \r\n"
|
||
" uint local_start_offset=kernel_index*N; \r\n"
|
||
" for(int i=0; i<N; i++) \r\n"
|
||
" { \r\n"
|
||
" data[i+local_start_offset] *= 2.0; \r\n"
|
||
" } \r\n"
|
||
" } \r\n";
|
||
//+------------------------------------------------------------------+
|
||
//| Test_CPU |
|
||
//+------------------------------------------------------------------+
|
||
bool Test_CPU(double &data[],const int N,const int id,const int total_arrays)
|
||
{
|
||
//--- check array size
|
||
if(ArraySize(data)==0) return(false);
|
||
//--- check array index
|
||
if(id>total_arrays) return(false);
|
||
//--- calculate local offset for array with index id
|
||
int local_start_offset=id*N;
|
||
//--- multiply elements by 2
|
||
for(int i=0; i<N; i++)
|
||
{
|
||
data[i+local_start_offset]*=2.0;
|
||
}
|
||
return true;
|
||
}
|
||
//---
|
||
#define ARRAY_SIZE 100 // size of the array
|
||
#define TOTAL_ARRAYS 5 // total arrays
|
||
//--- OpenCL handles
|
||
int cl_ctx; // OpenCL context handle
|
||
int cl_prg; // OpenCL program handle
|
||
int cl_krn; // OpenCL kernel handle
|
||
int cl_mem; // OpenCL buffer handle
|
||
//---
|
||
double DataArray1[]; // data array for CPU calculation
|
||
double DataArray2[]; // data array for GPU calculation
|
||
//+------------------------------------------------------------------+
|
||
//| Script program start function |
|
||
//+------------------------------------------------------------------+
|
||
int OnStart()
|
||
{
|
||
//--- initialize OpenCL objects
|
||
//--- create OpenCL context
|
||
if((cl_ctx=CLContextCreate())==INVALID_HANDLE)
|
||
{
|
||
Print("OpenCL not found. Error=",GetLastError());
|
||
return(1);
|
||
}
|
||
//--- create OpenCL program
|
||
if((cl_prg=CLProgramCreate(cl_ctx,cl_src))==INVALID_HANDLE)
|
||
{
|
||
CLContextFree(cl_ctx);
|
||
Print("OpenCL program create failed. Error=",GetLastError());
|
||
return(1);
|
||
}
|
||
//--- create OpenCL kernel
|
||
if((cl_krn=CLKernelCreate(cl_prg,"Test_GPU"))==INVALID_HANDLE)
|
||
{
|
||
CLProgramFree(cl_prg);
|
||
CLContextFree(cl_ctx);
|
||
Print("OpenCL kernel create failed. Error=",GetLastError());
|
||
return(1);
|
||
}
|
||
//--- create OpenCL buffer
|
||
if((cl_mem=CLBufferCreate(cl_ctx,ARRAY_SIZE*TOTAL_ARRAYS*sizeof(double),CL_MEM_READ_WRITE))==INVALID_HANDLE)
|
||
{
|
||
CLKernelFree(cl_krn);
|
||
CLProgramFree(cl_prg);
|
||
CLContextFree(cl_ctx);
|
||
Print("OpenCL buffer create failed. Error=",GetLastError());
|
||
return(1);
|
||
}
|
||
//--- set OpenCL kernel constant parameters
|
||
CLSetKernelArgMem(cl_krn,0,cl_mem);
|
||
CLSetKernelArg(cl_krn,1,ARRAY_SIZE);
|
||
CLSetKernelArg(cl_krn,2,TOTAL_ARRAYS);
|
||
//--- prepare data arrays
|
||
ArrayResize(DataArray1,ARRAY_SIZE*TOTAL_ARRAYS);
|
||
ArrayResize(DataArray2,ARRAY_SIZE*TOTAL_ARRAYS);
|
||
//--- fill arrays with data
|
||
for(int j=0; j<TOTAL_ARRAYS; j++)
|
||
{
|
||
//--- calculate local start offset for jth array
|
||
uint local_offset=j*ARRAY_SIZE;
|
||
//--- prepare array with index j
|
||
for(int i=0; i<ARRAY_SIZE; i++)
|
||
{
|
||
//--- fill arrays with function MathCos(i+j);
|
||
DataArray1[i+local_offset]=MathCos(i+j);
|
||
DataArray2[i+local_offset]=MathCos(i+j);
|
||
}
|
||
};
|
||
//--- test CPU calculation
|
||
for(int j=0; j<TOTAL_ARRAYS; j++)
|
||
{
|
||
//--- calculation of the array with index j
|
||
Test_CPU(DataArray1,ARRAY_SIZE,j,TOTAL_ARRAYS);
|
||
}
|
||
//--- prepare CLExecute params
|
||
uint offset[]={0};
|
||
//--- global work size
|
||
uint work[]={TOTAL_ARRAYS};
|
||
//--- write data to OpenCL buffer
|
||
CLBufferWrite(cl_mem,DataArray2);
|
||
//--- execute OpenCL kernel
|
||
CLExecute(cl_krn,1,offset,work);
|
||
//--- read data from OpenCL buffer
|
||
CLBufferRead(cl_mem,DataArray2);
|
||
//--- total error
|
||
double total_error=0;
|
||
//--- compare results and calculate error
|
||
for(int j=0; j<TOTAL_ARRAYS; j++)
|
||
{
|
||
//--- calculate local offset for jth array
|
||
uint local_offset=j*ARRAY_SIZE;
|
||
//--- compare the results
|
||
for(int i=0; i<ARRAY_SIZE; i++)
|
||
{
|
||
double v1=DataArray1[i+local_offset];
|
||
double v2=DataArray2[i+local_offset];
|
||
double delta=MathAbs(v2-v1);
|
||
total_error+=delta;
|
||
//--- show first and last arrays
|
||
if((j==0) || (j==TOTAL_ARRAYS-1))
|
||
PrintFormat("array %d of %d, element [%d]: %f, %f, [error]=%f",j+1,TOTAL_ARRAYS,i,v1,v2,delta);
|
||
}
|
||
}
|
||
PrintFormat("Total error: %f",total_error);
|
||
//--- delete OpenCL objects
|
||
//--- free OpenCL buffer
|
||
CLBufferFree(cl_mem);
|
||
//--- free OpenCL kernel
|
||
CLKernelFree(cl_krn);
|
||
//--- free OpenCL program
|
||
CLProgramFree(cl_prg);
|
||
//--- free OpenCL context
|
||
CLContextFree(cl_ctx);
|
||
//---
|
||
return(0);
|
||
}
|
||
|
||
```
|