Module tornado.api
Package uk.ac.manchester.tornado.api

Class KernelContext

java.lang.Object
uk.ac.manchester.tornado.api.KernelContext
All Implemented Interfaces:
ExecutionContext
public class KernelContext extends Object implements ExecutionContext
Context of TornadoVM execution to exploit kernel-parallel applications, in which the parallelism is implicit.

The application can access thread-id for 1D, 2D and 3D dimensions. Additionally, the application can access local memory (OpenCL terminology), or shared memory (CUDA terminology) as well as synchronization primitives such as barriers.

  • KernelContext is an object exposed by the TornadoVM API in order to leverage low-level programming features provided by heterogeneous frameworks (e.g. OpenCL, CUDA) to the developers, such as thread-id, access to local/shared memory and barriers.
  • KernelContext provides a Java API that is transparently translated to both OpenCL and PTX by the TornadoVM JIT compiler. The main difference with the TaskGraph API is that the tasks within a TaskGraph that use KernelContext must be GridScheduler.

  • Field Summary

    Fields
    Modifier and Type
    Field
    Description
    final Integer
    globalGroupSizeX
    It returns the global group size of a particular dimension (e.g.
    final Integer
    globalGroupSizeY
     
    final Integer
    globalGroupSizeZ
     
    final Integer
    globalIdx
    It returns the thread identifier for the first dimension.
    final Integer
    globalIdy
    It returns the thread identifier for the second dimension.
    final Integer
    globalIdz
    It returns the thread identifier for the third dimension.
    final Integer
    groupIdx
     
    final Integer
    groupIdy
     
    final Integer
    groupIdz
     
    final Integer
    localGroupSizeX
    It returns the global group size of a particular dimension (e.g.
    final Integer
    localGroupSizeY
     
    final Integer
    localGroupSizeZ
     
    final Integer
    localIdx
     
    final Integer
    localIdy
     
    final Integer
    localIdz
     

  • Constructor Summary

    Constructors
    Constructor
    Description
    KernelContext()
    Class constructor specifying a particular WorkerGrid object.

  • Method Summary

    Modifier and Type
    Method
    Description
    byte[]
    allocateByteLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    double[]
    allocateDoubleLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    float[]
    allocateFloatLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    HalfFloat[]
    allocateHalfFloatLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    int[]
    allocateIntLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    long[]
    allocateLongLocalArray(int size)
    It allocates a single dimensional array in local memory (known as shared memory in PTX).
    void
    atomicAdd(int[] array, int index, int val)
    Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.
    void
    atomicAdd(DoubleArray array, int index, double val)
    Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.
    void
    atomicAdd(FloatArray array, int index, float val)
    Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.
    void
    atomicAdd(IntArray array, int index, int val)
    Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.
    void
    atomicAdd(LongArray array, int index, long val)
    Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.
    void
    globalBarrier()
    Method used as a barrier to synchronize the order of memory operations to the global memory.
    void
    localBarrier()
    Method used as a barrier to synchronize the order of memory operations to the local memory (known as shared memory in PTX).

    Methods inherited from class java.lang.Object

    clone, equals, finalize, getClass, hashCode, notify, notifyAll, toString, wait, wait, wait

  • Field Details

    • globalIdx

      public final Integer globalIdx
      It returns the thread identifier for the first dimension.

      OpenCL equivalent: get_global_id(0);

      PTX equivalent: blockIdx.x * blockDim.x + threadIdx.x

    • globalIdy

      public final Integer globalIdy
      It returns the thread identifier for the second dimension.

      OpenCL equivalent: get_global_id(1);

      PTX equivalent: blockIdx.y * blockDim.y + threadIdx.y

    • globalIdz

      public final Integer globalIdz
      It returns the thread identifier for the third dimension.

      OpenCL equivalent: get_global_id(2);

      PTX equivalent: blockIdx.z * blockDim.z + threadIdx.z

    • groupIdx

      public final Integer groupIdx

    • groupIdy

      public final Integer groupIdy

    • groupIdz

      public final Integer groupIdz

    • localIdx

      public final Integer localIdx

    • localIdy

      public final Integer localIdy

    • localIdz

      public final Integer localIdz

    • globalGroupSizeX

      public final Integer globalGroupSizeX
      It returns the global group size of a particular dimension (e.g. X, Y, Z).

      OpenCL equivalent: get_global_size();

      PTX equivalent: gridDim * blockDim

    • globalGroupSizeY

      public final Integer globalGroupSizeY

    • globalGroupSizeZ

      public final Integer globalGroupSizeZ

    • localGroupSizeX

      public final Integer localGroupSizeX
      It returns the global group size of a particular dimension (e.g. X, Y, Z).

      OpenCL equivalent: get_local_size();

      PTX equivalent: blockDim

    • localGroupSizeY

      public final Integer localGroupSizeY

    • localGroupSizeZ

      public final Integer localGroupSizeZ

  • Constructor Details

    • KernelContext

      public KernelContext()
      Class constructor specifying a particular WorkerGrid object.

  • Method Details

    • localBarrier

      public void localBarrier()
      Method used as a barrier to synchronize the order of memory operations to the local memory (known as shared memory in PTX).

      OpenCL equivalent: barrier(CLK_LOCAL_MEM_FENCE);

      PTX equivalent: barrier.sync;

      Specified by:
      localBarrier in interface ExecutionContext

    • globalBarrier

      public void globalBarrier()
      Method used as a barrier to synchronize the order of memory operations to the global memory.

      OpenCL equivalent: barrier(CLK_GLOBAL_MEM_FENCE);

      PTX equivalent: barrier.sync;

      Specified by:
      globalBarrier in interface ExecutionContext

    • allocateIntLocalArray

      public int[] allocateIntLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateIntLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      int[]: reference to the int array

    • allocateByteLocalArray

      public byte[] allocateByteLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateByteLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      byte[]: reference to the byte array

    • allocateHalfFloatLocalArray

      public HalfFloat[] allocateHalfFloatLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateHalfFloatLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      HalfFloatArray: reference to the byte array

    • allocateLongLocalArray

      public long[] allocateLongLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateLongLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      long[]: reference to the int array

    • allocateFloatLocalArray

      public float[] allocateFloatLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateFloatLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      float[]: reference to the int array

    • allocateDoubleLocalArray

      public double[] allocateDoubleLocalArray(int size)
      It allocates a single dimensional array in local memory (known as shared memory in PTX).
      Specified by:
      allocateDoubleLocalArray in interface ExecutionContext
      Parameters:
      size - the size of the array
      Returns:
      double[]: reference to the int array

    • atomicAdd

      public void atomicAdd(IntArray array, int index, int val)
      Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.

      PTX equivalent: atomicAdd(int* address, int val);

      Specified by:
      atomicAdd in interface ExecutionContext

    • atomicAdd

      public void atomicAdd(int[] array, int index, int val)
      Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.

      PTX equivalent: atomicAdd(int* address, int val);

      Specified by:
      atomicAdd in interface ExecutionContext

    • atomicAdd

      public void atomicAdd(LongArray array, int index, long val)
      Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.

      PTX equivalent: atomicAdd(long* address, long val);

      Specified by:
      atomicAdd in interface ExecutionContext

    • atomicAdd

      public void atomicAdd(FloatArray array, int index, float val)
      Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.

      PTX equivalent: atomicAdd(float* address, float val);

      Specified by:
      atomicAdd in interface ExecutionContext

    • atomicAdd

      public void atomicAdd(DoubleArray array, int index, double val)
      Method used to read a memory address by using the array and the index, then add the value of val to it, and write the result back to the same address.

      PTX equivalent: atomicAdd(double* address, double val);

      Specified by:
      atomicAdd in interface ExecutionContext