search

hughperkins / cltorch

An OpenCL backend for torch.
Other
289 stars 26 forks source link

readme

cltorch

An OpenCL backend for torch.

What is this?

It's a high-performance matrix library for OpenCL, that runs on your GPU(s) harnessing the massive computational capacity that they provide.

Most of the standard operations in torch are supported. If there are any missing that you need, please raise an issue.

What's working

Most things really :-) Detailed description at ImplementedDetails.md. Please don't hesitate to raise an issue for anything that's missing that you would like to see added.

Installation

IMPORTANT! THIS HAS CHANGED. Please install a specific Torch distro, as described below. Simply doing luarocks install cltorch is no longer supported

Please see distro-cl for installation instructions.

cltorch-specific features

The following features are either cltorch-specific, or do not exist in cutorch:

feature in torch? in cutorch? in cltorch?
apply/map/map2 Yes Via an extension Yes
profiling tools Via CUDA nvvp Yes
point tensors Yes
custom user kernels Not applicable Via cutorch-rtc Yes

apply/map/map2

apply, map, map2 exist in torch, but how to make them work on the GPU? Cannot just pass in lua functions.

What we do is, you can provide opencl code directly to apply_on_gpu, map_on_gpu and map2_on_gpu, as a string expression. This will run on the gpu, at full speed. Examples, for x, y, z being identically sized torch.ClTensors:

x:apply_on_gpu("x = sqrt(x + 3.5)")
x:map_on_gpu(y, "x = 1000 * x + y * 10")
x:map2_on_gpu(y, z, "x = sqrt(1000 * x + y * 10 + z * z)")

Profiling tools

Following tools are available to aid with profiling:

Method Description
cltorch.setProfiling(1) turn on opencl kernel profiling
cltorch.dumpProfiling() dump opencl kernel profiling timings since last call
cltorch.setEnableTiming(1) enable collection of cumulative wall-clock timings for cltorch code
cltorch.dumpTimings() dump cumulative wall-clock timings for cltorch code
cltorch.setTrace(1) print all gpu buffer allocations and copies between host/gpu

OpenCL Profiling

OpenCL natively provides facilities to measure the execution time of kernels, without needing to call cltorch.finish() or similar first, using clGetEventProfilingInfo. In cltorch, you dont need to know how this works ;-) Simply call, at the start of your code:

cltorch.setProfiling(1)

Then, after running the piece of code under scrutiny, simply call:

cltorch.dumpProfiling()

Timings are cumulative across multiple calls to the same kernel.

DumpTimings

This uses the wall-clock times to measure the elapsed time in different sections of cltorch code. The way it works is, each time the cltorch c++ code calls StatefulTimer::instance()->timeCheck("some status"), the wall-clock time since the last call to ->timeCheck() will be added to the cumulative time for some status. You can pass any status as a string. Then, after running the piece of code under the scrutiny, in your Lua program, simply call cltorch.dumpTimings() to dump these cumulative timings.

Update: please first call cltorch.setEnableTiming(true) to enable collection of timing information. This is global across all devices.

GPU buffer allocations and copies

You can log all GPU buffer allocations, copies to host, and copies to GPU device. Simply call:

cltorch.setTrace(1)

Any buffer allocations, and copies between host and device, will now be printed to stdout.

Point tensors: reduce pipeline stalls

Point tensors help to eliminate pipeline stalls associated with ReduceAll operations such as sometensor:sum(). Why does :sum() cause pipeline stalls, and how do point tensors eliminate this source of stalls?

If we send a single instruction (a kernel) to the gpu, there will be some latency whilst the instruction arrives at the gpu, and starts running, and some more latency after the calculations have finished, whilst the results are retrieved back from the GPU. Maybe we send:

a:add(1)

We can draw a picture of what happens. Time is towards the right. GPU is at the top. CPU at the bottom:

gpu single instruction

But we can send lots of instructions, without waiting for the earlier ones to finish. Maybe we do:

a:add(b)
a:mul(3)
b:mul(a)
c:add(a)

This might look like this, we dont have to wait for the previous instruction to finish: gpu pipeline

But now imagine what happens if we process the following instruction:

a:div(a:sum())

Looks like this: gpu stall

Classic reduceall => Massive pipeline stall

Point tensors eliminate this. When we do the reduceall, the :sum() operation, we keep the results on the gpu, like this:

c = torch.Tensor(20,30):uniform():cl() -- create a tensor on the GPU
res = torch.ClTensor()                 -- create a point tensor on the GPU
res:sum(c)                             -- sum c, and keep the result in res, on the GPU

res is a point tensor. It has zero dimensions. It contains a single scalar float. It stays on the GPU. We can feed it into other operations as follows:

c:div(res)  -- divide c by res

We can send this instruction straight away, even before the first :sum(c) instruction has arrived at the GPU. So, no more stall.

By the way, it's possible to print the value of a point tensor, by printing it, or calling the :s() operator. Normally you wouldnt do this except during debugging though, since obviously this will need to wait for the gpu operation to finish, and for the data to come all the way back from the GPU :-)

Custom user kernels

Custom user kernels let you run OpenCL code directly from Torch Lua! Of course, you can already do this with apply, map, and map2, see above. But now you can provide whole kernel functions, and other functions, and pass ClTensors into these kernels!

Example of how to use:

require 'cltorch'

k = torch.ClKernel({input={nElements='int', input='torch.ClTensor'},output={output='torch.ClTensor'},src=[[
   int linearId = get_global_id(0);
   if(linearId < nElements) {
     output_data[linearId] = input_data[linearId] + 3.0f;
   }
]]})
print('k', k)
k:print()

x = torch.ClTensor({3,5,2})
y = torch.ClTensor({6,4,2})
print('x before\n', x)
print('y before\n', y)

k:run({nElements=3, input=x, output=y})

print('y after\n', y)

Output from this example:

Using Intel platform: Intel Gen OCL Driver
Using device: Intel(R) HD Graphics BroadWell U-Processor GT2
k   torch.ClKernel
Original source
===============
   int linearId = get_global_id(0);
   if(linearId < nElements) {
     output_data[linearId] = input_data[linearId] + 3.0f;
   }

Generated source
================
typedef struct THClTensorInfoCl {
  unsigned int sizes[25];
  unsigned int strides[25];
  int offset;
  int dims;
} TensorInfoCl;

kernel void user_kernel(
    global struct THClTensorInfoCl *input_info, global float * input_data,
    int nElements,
    global struct THClTensorInfoCl *output_info, global float * output_data
) {
   int linearId = get_global_id(0);
   if(linearId < nElements) {
     output_data[linearId] = input_data[linearId] + 3.0f;
   }

}

x before
     3
 5
 2
[torch.ClTensor of size 3]

y before
     6
 4
 2
[torch.ClTensor of size 3]

y after
     6
 8
 5
[torch.ClTensor of size 3]

If you want, you can specify the number of workgroups, and the workgroupsize, yourself:

k:run({nElements=3, input=x, output=y}, {numWorkgroups=10, workgroupSize=32}

Co-existence with cutorch

Third-party libraries

cltorch uses the following libraries. These are automatically built as part of cltorch build process:

At runtime, if you want to call any of the cltorch methods, you will also need:

Guidelines for contributors

You might or might not find ContributorGuidelines.md useful. Not required reading, but it is there if you want to see my own thoughts and ideas on how I am currently approaching cltorch development, and cutorch-porting.

Also, some more technical guidelines on porting, in the clnn repository, at porting-guidelines.md.

Related projects

There is an OpenCL backend for nn and nngraph at clnn.

There is an HCC backend for Torch at: hctorch

Recent changes

OlderChanges.md