OpenCL

OpenCL (engl. Open Computing Language) on Khronos Groupin kehittämä rinnakkaisohjelmointiin tarkoitettu viitekehys, jonka tarkoitus on mahdollistaa ohjelmistojen hajauttaminen erilaisille laskenta-alustoille verrattain helposti. 
OpenCL koostuu kohdelaitteilla suoritettavista ydinfunktioista (kernel) sekä isäntälaitteella suoritettavista ohjauskäskyistä. Isäntälaitteet ovat tyypillisesti tavanomaisia CPU-laitteita, joissa suoritettava koodi voidaan kirjoittaa eri ohjelmointikielillä, kunhan ne tukevat OpenCL APIa (esimerkiksi C/C++, C#, Java, Fortran, Python, ...). OpenCL on pyritty suunnittelemaan siten, että kohdealustana voivat toimia hyvinkin erilaiset suoritusympäristöt, kuten CPU, GPU, DSP,... Rinnakkaistaminen on mahdollista toteuttaa myös useammalla alustalla yhtä aikaa (esim NxCPU + MxGPU). OpenCL skaalautuu aina sulautetuista alustoista supertietokoneisiin. Suurimpia laitevalmistajia jotka tukevat OpenCL-rajapintaa ovat esimerkiksi NVIDIA, AMD, Apple ja Intel. 
Kernel-sovellukset kehitetään OpenCL C -kielellä, joka perustuu C99-kieleen ja sisältää laajennuksia, jotka helpottavat rinnakkaissuorituksen hallittavuutta. Tällaisia laajennuksia ovat esimerkiksi osoiteavaruuden indeksointi globaalisti sekä lokaalisti rinnakkaistetussa tehtävälohkossa. 
OpenCL-alustoista huomattava osa koostuu erilaisista grafiikkapiireistä ja siksi on kieleen lisätty myös liittymärajapinnat sekä OpenGL- että DirectX-muistiobjekteihin. Käytännössä tämä näkyy siten, että esimerkiksi OpenGL-tekstuureita voidaan käsitellä suoraan OpenCL-kerneleillä.
OpenCL mahdollistaa sekä dataperusteisen rinnakkaistamisen että tehtävälähtöisen rinnakkaistamisen. Datalähtöisessä ohjelmointimallissa kernel-funktiot suorittavat samaa käskysarjaa eri tietolohkoilla yhtä aikaa (SIMD). Tehtävälähtöisessä tapauksessa työ voidaan jakaa eri kerneleihin, joiden työryhmiä suoritetaan yhtä aikaa (task parallel). Kernel-työryhmien suoritusta voidaan synkronoida tehtäväjonoilla ja näihin syötettävillä synkronointikäskyillä.

SPIR ja SPIR-V

Kohdelaitteiden suuren määrän ja vaihtelevuuden vuoksi OpenCL -ohjelmakoodi käännetään välimuotoiselle SPIR (engl. Standard Portable Intermediate Representation) muodolle. SPIR käännetään kohdelaitteella alustakohtaiselle käskykannalle ja arkkitehtuurille. SPIR voidaan verrata myös HSAIL-käännökseen.
SPIR-V on avoin alustariippumaton standardi.

Standardiversiot

Ensimmäinen versio OpenCL 1.0 julkaistiin joulukuussa 2008. 
Kesäkuussa 2010 esiteltiin laajennus 1.1, jota useat laitevalmistajat tukevat. 
Spesifikaatiot laajennukseen OpenCL 1.2 esiteltiin marraskuussa 2011.
Viimeisin julkaistu versio tammikuussa 2017 on OpenCL 2.1.

Esimerkki

Tämä esimerkki toteuttaa Fourier’n muunnos (FFT).
// create a compute context with GPU devicecontext = clCreateContextFromType(CL_DEVICE_TYPE_GPU);// create a work-queuequeue = clCreateWorkQueue(context, NULL, NULL, 0);// allocate the buffer memory objectsmemobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA);memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL);// create the compute programprogram = clCreateProgramFromSource(context, 1, &fft1D_1024_kernel_src, NULL);// build the compute program executableclBuildProgramExecutable(program, false, NULL, NULL);// create the compute kernelkernel = clCreateKernel(program, “fft1D_1024”);// create N-D range object with work-item dimensionsglobal_work_size[0] = n;local_work_size[0] = 64;range = clCreateNDRangeContainer(context, 0, 1, global_work_size, local_work_size);// set the args valuesclSetKernelArg(kernel, 0, (void *)&memobjs[0], sizeof(cl_mem), NULL);clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem), NULL);clSetKernelArg(kernel, 2, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);clSetKernelArg(kernel, 3, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);

 // execute kernelclExecuteKernel(queue, kernel, NULL, range, NULL, 0, NULL);
Todellinen laskelma:
// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into // calls to a radix 16 function, another radix 16 function and then a radix 4 function __kernel void fft1D_1024 (__global float2 *in, __global float2 *out, 
                          __local float *sMemx, __local float *sMemy) { 
  int tid = get_local_id(0); 
  int blockIdx = get_group_id(0) * 1024 + tid; 
  float2 data[16]; 
  // starting index of data to/from global memory 
  in = in + blockIdx;  out = out + blockIdx; 
  globalLoads(data, in, 64); // coalesced global reads 
  fftRadix16Pass(data);      // in-place radix-16 pass 
  twiddleFactorMul(data, tid, 1024, 0); 
  // local shuffle using local memory 
  localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4))); 
  fftRadix16Pass(data);               // in-place radix-16 pass 
  twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication 
  localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15))); 
  // four radix-4 function calls 
  fftRadix4Pass(data); fftRadix4Pass(data + 4); 
  fftRadix4Pass(data + 8); fftRadix4Pass(data + 12); 
  // coalesced global writes 
  globalStores(data, out, 64); 
}