Device

The device abstracts physical devices by its responsibilities

  • Memory management
  • Building kernels (device functions)
  • Managing streams of execution

Constructors/Destructors

Constructors

occa::device::device (
const occa::properties &props
)

Creates a occa::device object from occa::properties

Parameters:

  • props  -  

    Properties used to initialize the device

occa::device device("mode: 'Serial'");
occa::properties props;
props["mode"] = "Serial";
occa::device device(props);
New in version 1.0
void occa::device::setup (
const occa::properties &props
)

Initializes the occa::device object from occa::properties

Parameters:

  • props  -  

    Properties used to initialize the device

occa::device device;
device.setup("mode: 'Serial'");
occa::device device;

occa::properties props;
props["mode"] = "Serial";

device.setup(props);
New in version 1.0

Destructors

void occa::device::free ()

Frees

  • The occa::device object
  • All occa::stream objects created by this device
  • The backend device handle
occa::device device("mode: 'Serial'");
device.free();
New in version 0.1

Device Information

bool occa::device::isInitialized ()

Returns true if the device has been initialized by

  • constructor
  • device::setup()
occa::device device;
std::cout << "Device initialized: " << device.isInitialized() << '\n';
device.setup("mode: 'Serial'");
std::cout << "Device initialized: " << device.isInitialized() << '\n';
New in version 1.0
const std::string& occa::device::mode ()

Returns the mode of the device. The mode can also be obtained through its properties.

occa::device device("mode: 'Serial'");
std::cout << "The device mode is: " << device.mode() << '\n';
New in version 0.1
const occa::properties& occa::device::properties ()

Returns the occa::properties used to initialize the device.

occa::device device("mode: 'Serial'");
std::cout << "The device mode is: " << device.properties()["mode"] << '\n';
New in version 1.0
occa::udim_t occa::device::memorySize ()

Returns the memory capacity in the device. This value is not decreased by allocating memory on the device.

occa::device device("mode: 'Serial'");
std::cout << "Memory in device: " << device.memorySize() << '\n';
New in version 0.1
occa::udim_t occa::device::memoryAllocated ()

Returns the memory allocated using this occa::device. Memory allocated on the same physical device through different occa::device objects or non-OCCA methods are not included.

occa::device device("mode: 'Serial'");
std::cout << "Memory allocated in device: " << device.memoryAllocated() << '\n';
New in version 0.1
void occa::device::hasSeparateMemorySpace ()

Returns whether the memory in the connected physical device is not located inside of the host memory.

occa::device device("mode: 'Serial'");
std::cout << "Device shared memory with host: "
          << device.hasSeparateMemorySpace() << '\n';
occa::device device("mode: 'CUDA', deviceID: 0");
std::cout << "Device shared memory with host: "
          << device.hasSeparateMemorySpace() << '\n';
New in version 1.0

Memory Allocation

occa::memory occa::device::malloc (
const dim_t bytes,
void *src = NULL,
const occa::properties &props = occa::properties()
)

Allocates bytes bytes in device memory. If src is passed, the newly allocated memory is initialized with data from src.

Some devices might have muliple forms of memory allocation. The optional props argument is present for exposing such features in future backends.

Parameters:

  • bytes  -  

    Bytes allocated for the returned occa::memory

  • src  -  

    If passed, data from src initializes the occa::memory

  • props  -  

    Additional information to initialize the occa::memory

occa::device device("mode: 'Serial'");
occa::memory array = device.malloc(5 * sizeof(int));
occa::device device("mode: 'Serial'");
int *array = new int[5];
for (int i = 0; i < 5; ++i) {
  array[i] = i;
}
occa::memory deviceArray = device.malloc(5 * sizeof(int), array);
New in version 0.1
void* occa::device::umalloc (
const dim_t bytes,
void *src = NULL,
const occa::properties &props = occa::properties()
)

Similar to device::malloc, device::umalloc allocates bytes bytes in device memory. However, the return value is a pointer in host memory of the same size. The host and device data are synchronized through device::finish() or other UVA methods.

If src is passed, the newly allocated memory is initialized with data from src.

Some devices might have muliple forms of memory allocation. The optional props argument is present for exposing such features in future backends.

Parameters:

  • bytes  -  

    Bytes allocated for the returned occa::memory

  • src  -  

    If passed, data from src initializes the occa::memory

  • props  -  

    Additional information to initialize the occa::memory

occa::device device("mode: 'Serial'");

// Host  : [?, ?, ?, ?, ?]
// Device: [?, ?, ?, ?, ?]
int *array = (int*) device.umalloc(5 * sizeof(int));

// Host  : [0, 1, 2, 3, 4]
// Device: [?, ?, ?, ?, ?]
for (int i = 0; i < 5; ++i) {
  array[i] = i;
}

// Host  : [0, 1, 2, 3, 4]
// Device: [0, 1, 2, 3, 4]
occa::syncToDevice(array);
New in version 1.0

Building Kernels

occa::kernel occa::device::buildKernel (
const std::string &filename,
const std::string &kernelName,
const occa::properties &props = occa::properties()
)

Builds and returns the kernel names kernelName in file filename. When props are passed, they can be used to pass additional information to the backend.

For more information, check out the kernel section

Parameters:

  • filename  -  

    The filename containing the kernel to be built

  • kernelName  -  

    The name of the kernel to be built

  • props  -  

    Additional information to build the occa::kernel

occa::device device("mode: 'Serial'");
occa::kernel addVectors = device.buildKernel("addVectors.okl",
                                             "addVectors");
New in version 0.1
occa::kernel occa::device::buildKernelFromString (
const std::string &content,
const std::string &kernelName,
const occa::properties &props = occa::properties()
)

Builds and returns the kernel names kernelName in source code content. When props are passed, they can be used to pass additional information to the backend.

For more information, check out the kernel section

Parameters:

  • content  -  

    Source code containing kernel to be built

  • kernelName  -  

    The name of the kernel to be built

  • props  -  

    Additional information to build the occa::kernel

occa::device device("mode: 'Serial'");
occa::kernel addVectors = device.buildKernelFromString(
  "@kernel void addVectors(const int entries,\n"
  "                        const float *a,\n"
  "                        const float *b,\n"
  "                        float *ab) {\n"
  "  for (int i = 0; i < entries; ++i; tile(16)) {\n"
  "    if (i < entries)\n"
  "      ab[i] = a[i] + b[i];\n"
  "  }\n"
  "}\n"
);
New in version 0.1
occa::kernel occa::device::buildKernelFromBinary (
const std::string &filename,
const std::string &kernelName,
const occa::properties &props = occa::properties()
)

Returns the already-compiled kernel kernelName from binary file filename. When props are passed, they can be used to pass additional information to the backend.

For more information, check out the kernel section

Parameters:

  • filename  -  

    The filename of the binary containing the already-compiled kernel

  • kernelName  -  

    The name of the kernel in the binary

  • props  -  

    Additional information for fetching the kernel in the binary

occa::device device("mode: 'Serial'");
occa::kernel addVectors = device.buildKernelFromBinary("addVectors.so",
                                                       "addVectors");
New in version 0.1

Stream Management

void occa::device::finish ()

Blocks until all device kernel calls and memory transfers complete.

occa::device device("mode: 'Serial'");
// Launching myKernel is non-blocking
myKernel(1);
// Waits until myKernel is done
device.finish();
occa::device device("mode: 'Serial'");
// Asynchronous copy
occa::memcpy(dest, src,
             10 * sizeof(int),
             "async: true");
// Waits until the async copy is done
device.finish();
New in version 0.1

Stream Creation

occa::stream occa::device::createStream ()

Creates and returns a new occa::stream.

occa::device device("mode: 'Serial'");
occa::stream stream1 = device.getStream();
occa::stream stream2 = device.createStream();
device.setStream(stream2);
New in version 0.1
occa::stream occa::device::getStream ()

Returns the occa::stream the device is currently using.

occa::device device("mode: 'Serial'");
occa::stream stream1 = device.getStream();
occa::stream stream2 = device.createStream();
device.setStream(stream2);
New in version 0.1
void occa::device::setStream (
occa::stream stream
)

Sets the stream the device will use for all following kernel calls and memory transfers.

Parameters:

  • stream  -  

    All following kernel calls and memory transfers are done using this stream

occa::device device("mode: 'Serial'");
occa::stream stream1 = device.getStream();
occa::stream stream2 = device.createStream();
device.setStream(stream2);
New in version 0.1

Stream Timings

occa::streamTag occa::device::tagStream ()

Tags the stream at its current state. The tag can later be used to compute time between itself and another tag.

occa::device device("mode: 'Serial'");
occa::streamTag start = device.tagStream();
// Launch a non-blocking kernel call
myKernel(1)
occa::streamTag end = device.tagStream();
std::cout << "Time taken: " << device.timeBetween(start, end) << '\n';
New in version 0.1
void occa::device::waitFor (
occa::streamTag tag
)

Blocks until all kernel calls and memory transfers before tag have finished

Parameters:

  • tag  -  

    Blocking tag

occa::device device("mode: 'Serial'");
myKernel(1)
occa::streamTag myKernelEnd = device.tagStream();
myOtherKernel(1)
device.waitFor(myKernelEnd);
New in version 0.1
double occa::device::timeBetween (
const streamTag &start,
const streamTag &end
)

Returns the time in seconds it took events between start to end tags to finish. Blocks until all kernel calls and memory transfers before end have finished.

Parameters:

  • start  -  

    occa::tag specifying when to start the timing

  • end  -  

    occa::tag specifying when to end the timing

occa::device device("mode: 'Serial'");
occa::streamTag start = device.tagStream();
// Launch a non-blocking kernel call
myKernel(1)
occa::streamTag end = device.tagStream();
std::cout << "Time taken: " << device.timeBetween(start, end) << '\n';
New in version 0.1

Interoperability

void* occa::device::getHandle (
const occa::properties &props = occa::properties()
)

Returns a handle the occa:device uses to communicate with its backend. The return format is dependent on the backend and props.

Parameters:

  • props  -  

    Additional information that might be needed to get specific handle objects from the backend

occa::device device("mode: 'CUDA', deviceID: 0");
int deviceID = *((CUdevice*) device.getHandle());
New in version 1.0
occa::memory occa::device::wrapMemory (
void *handle,
const dim_t bytes,
const occa::properties &props = occa::properties()
)

Returns an occa::memory that uses the passed handle to communicate with the backend. The handle format is dependent on the backend and props.

Parameters:

  • handle  -  

    Representation of the memory object in the device backend

  • bytes  -  

    Size of the memory object being wrapped

  • props  -  

    Additional information that might be needed to wrap the occa::memory object

int *a = new int[10];

occa::device device("mode: 'Serial'");
occa::memory o_a = device.wrapMemory(a, 10 * sizeof(int));
New in version 0.1
occa::stream occa::device::wrapStream (
void *handle,
const occa::properties &props = occa::properties()
)

Returns an occa::stream that uses the passed handle to communicate with the backend. The handle format is dependent on the backend and props.

Parameters:

  • handle  -  

    Representation of the stream object in the device backend

  • props  -  

    Additional information that might be needed to wrap the occa::memory object

occa::device device("mode: 'CUDA', deviceID: 0");

cudaStream_t stream1;
cudaStreamCreate(&stream1);

occa::stream o_defaultStream = device.wrapStream(NULL);
occa::stream o_stream1       = device.wrapStream(&stream1);
New in version 0.1