Cloo performance? [ OpenCL CPU vs Pure .NET ]

I'm posting some numbers that might make you happy. This test intersects a ray with a number of triangles and it was run in both OpenCL (on GF 9600 GT) and .NET (release config).

10 triangles:

Cloo ticks: 952072,             milliseconds: 0
.NET ticks: 5891232,            milliseconds: 2

100 triangles

Cloo ticks: 950576,             milliseconds: 0
.NET ticks: 6017176,            milliseconds: 2

1 000 triangles

Cloo ticks: 1767872,            milliseconds: 0
.NET ticks: 7309984,            milliseconds: 2

10 000 triangles

Cloo ticks: 2695160,            milliseconds: 1
.NET ticks: 19414624,           milliseconds: 7

65 535 triangles

Cloo ticks: 9008704,            milliseconds: 3
.NET ticks: 107045968,          milliseconds: 40

The test will be included in the next release of Cloo.

Hi nythrix,

Can you provide the code you use for your performance test ? I would like to do some tests too .

Thx

It sounds great...

I have run it and I got the following error :

"InvalidKernelNameComputeException"

Here is the stack trace :

Cloo.dll!Cloo.ComputeException.ThrowIfError(OpenTK.Compute.CL10.ErrorCode errorCode = InvalidKernelName) Line 150 C#
Cloo.dll!Cloo.ComputeKernel.ComputeKernel(string functionName = "intersect", Cloo.ComputeProgram program = {ComputeProgram(92622517)}) Line 87 + 0x8 bytes C#
Cloo.dll!Cloo.ComputeProgram.CreateKernel(string functionName = "intersect") Line 230 + 0x1b bytes C#
ClooTester.exe!ClooTester.TriangleIntersector.Run() Line 97 + 0x10 bytes C#
ClooTester.exe!ClooTester.Program.Main() Line 20 + 0x1d bytes C#

I run the ATI SDK... with :
- Intel(R) Core(TM)2 Duo CPU E7400 @ 2.80GHz
- Graphic card : Intel Eaglelake

Strange, this doesn't occur with nVidia drivers. I have to find out whether this is caused by the kernel or by the ATI drivers. Anyway, I can't tell for sure until I get home tonight.

I get the same error on nVidia drivers:

------------------| Start OpenCL platform info |------------------
name: NVIDIA CUDA
version: OpenCL 1.0 CUDA 3.0.1
profile: FULL_PROFILE
vendor: NVIDIA Corporation
extensions:
+ cl_khr_byte_addressable_store
+ cl_khr_gl_sharing
+ cl_nv_compiler_options
+ cl_nv_device_attribute_query

devices:
name: GeForce 9600 GT
driver: 195.62
vendor: NVIDIA Corporation
extensions:
+ cl_khr_byte_addressable_store
+ cl_khr_gl_sharing
+ cl_nv_compiler_options
+ cl_nv_device_attribute_query
+
+ cl_khr_global_int32_base_atomics
+ cl_khr_global_int32_extended_atomics
-------------------| End OpenCL platform info |-------------------

------------------| Start Program binary |------------------
//
// Generated by NVIDIA NVPTX Backend for LLVM
//

.version 1.5
.target sm_11, texmode_independent, map_f64_to_f32

// Global Launch Offsets
.const[0] .s32 %_global_block_offset[3];
.const[0] .s32 %_global_launch_offset[3];
.const[0] .s32 %_global_num_groups[3];
.const[0] .s32 %_global_size[3];
.const[0] .u32 %_work_dim;

.const .align 8 .b8 def___internal_i2opi_d[144] = { 0x08, 0x5D, 0x8D, 0x1F, 0xB1, 0x5F, 0xFB, 0x6B, 0xEA, 0x92, 0x52, 0x8A, 0xF7, 0x39, 0x07, 0x3D, 0x7B, 0xF1, 0xE5, 0xEB, 0xC7, 0xBA, 0x27, 0x75, 0x2D, 0xEA, 0x5F, 0x9E, 0x66, 0x3F, 0x46, 0x4F, 0xB7, 0x09, 0xCB, 0x27, 0xCF, 0x7E, 0x36, 0x6D, 0x1F, 0x6D, 0x0A, 0x5A, 0x8B, 0x11, 0x2F, 0xEF, 0x0F, 0x98, 0x05, 0xDE, 0xFF, 0x97, 0xF8, 0x1F, 0x3B, 0x28, 0xF9, 0xBD, 0x8B, 0x5F, 0x84, 0x9C, 0xF4, 0x39, 0x53, 0x83, 0x39, 0xD6, 0x91, 0x39, 0x41, 0x7E, 0x5F, 0xB4, 0x26, 0x70, 0x9C, 0xE9, 0x84, 0x44, 0xBB, 0x2E, 0xF5, 0x35, 0x82, 0xE8, 0x3E, 0xA7, 0x29, 0xB1, 0x1C, 0xEB, 0x1D, 0xFE, 0x1C, 0x92, 0xD1, 0x09, 0xEA, 0x2E, 0x49, 0x06, 0xE0, 0xD2, 0x4D, 0x42, 0x3A, 0x6E, 0x24, 0xB7, 0x61, 0xC5, 0xBB, 0xDE, 0xAB, 0x63, 0x51, 0xFE, 0x41, 0x90, 0x43, 0x3C, 0x99, 0x95, 0x62, 0xDB, 0xC0, 0xDD, 0x34, 0xF5, 0xD1, 0x57, 0x27, 0xFC, 0x29, 0x15, 0x44, 0x4E, 0x6E, 0x83, 0xF9, 0xA2 };

.entry vectorAdd
(
.param .b32 vectorAdd_param_0,
.param .b32 vectorAdd_param_1,
.param .b32 vectorAdd_param_2
)
{
.reg .f32 %f<4>;
.reg .s32 %r<9>;

_vectorAdd:
{
// get_global_id(0)
.reg .u32 %vntidx;
.reg .u32 %vctaidx;
.reg .u32 %vtidx;
cvt.u32.u16 %vntidx, %ntid.x;
cvt.u32.u16 %vctaidx, %ctaid.x;
cvt.u32.u16 %vtidx, %tid.x;
mad.lo.s32 %r1, %vntidx, %vctaidx, %vtidx;
.reg .u32 %temp;
ld.const.u32 %temp, [%_global_launch_offset+0];
add.u32 %r1, %r1, %temp;
}

shl.b32 %r2, %r1, 2;
ld.param.u32 %r3, [vectorAdd_param_1];
ld.param.u32 %r4, [vectorAdd_param_0];
add.s32 %r5, %r4, %r2;
add.s32 %r6, %r3, %r2;
ld.param.u32 %r7, [vectorAdd_param_2];
ld.global.f32 %f1, [%r5];
ld.global.f32 %f2, [%r6];
add.rn.ftz.f32 %f3, %f1, %f2;
add.s32 %r8, %r7, %r2;
st.global.f32 [%r8], %f3;
ret;
}

-------------------| End Program binary |-------------------

------------------| Start Memory Object mapping |------------------
Original content:
2934802780230322176
8527569823997102080
204372612966262688
7046751857986121728
3974656866072502784
1309497746332879616
Mapped content:
2934802780230322176
8527569823997102080
204372612966262688
7046751857986121728
3974656866072502784
1309497746332879616
-------------------| End Memory Object mapping |-------------------

------------------| Start VectorAdd |------------------
56.57155 + 58.45175 = 115.0233
58.53479 + 70.96561 = 129.5004
7.859627 + 63.40628 = 71.26591
6.03061 + 52.21336 = 58.24397
49.79776 + 95.02283 = 144.8206
46.15215 + 2.900972 = 49.05312
57.18979 + 83.5145 = 140.7043
72.97466 + 0.7222652 = 73.69693
8.511994 + 2.004705 = 10.5167
75.73506 + 20.94896 = 96.68402
-------------------| End VectorAdd |-------------------

------------------| Start Triangle intersection |------------------
:3: error: invalid address space for argument to __kernel function
global float4 dir,

So,

I have do some test and in my "Test" application I receive the following message "Compilation failed".

But when I put my "CL" code in your application I receive a complete error description by calling :
program.GetBuildLog(context.Devices[0])

Here is the code I use :
ComputeContext context = new ComputeContext(DeviceTypeFlags.DeviceTypeDefault, null, null);
ComputeProgram program = new ComputeProgram(context, kernelSource);

try
{
program.Build(context.Devices, null, null, IntPtr.Zero);
}
catch (ComputeException e)
{
LogManager.Trace(LogSeverity.CriticalError, LogType.Rendering, program.GetBuildLog(context.Devices[0]));
return;
}

Why I receive a "short" description of the error and not the complete one like in your application ?

Thx

Since this topic has grown a bit off-topic, please direct further posts to the bug report I've created. Thanks.

PS.: I will probably not fix this issue before weekend. I'm out of time for the next two days, guys.

Hi,

1 - I think that there is an error in the way we compute the "Performance". When using OpenCL we must :
a) take account of the "Data transfer" between .NET memory and the CPU or GPU memory (Theses are 2 kinds of different transfers.
b) we must take into account the fact that the code is compiled... in some case we have to create a new context, a new compilation and a new data transfer
c) we need to free some OpenCL memory, in some cases too

2 - Using the .NET Parallel LINQ will not speed up the processing up to 50 times... except if you have at least 50 core on your CPU....

So, it is very difficult to compare both !

What I suggest is that we should put the 'CL data creation + CL data transfer' in the speed test.

Regards

Hi Anto,

I understand your point of view and agree...

Except that in some cases it is not possible to directly convert .NET data to OpenCL data, in fact, we can do some "basic performance tests" but in the reality there are "some cases" where we must take "more" parameters into account to compute the "real" performance gain ;-)

Even, in some case we can send all the "triangles" information at the startup of the application and doing all the "intersections tests" without doing the "transfer" (Set*Arg) at each call...

When playing with this kind of technology we must "adapt" the "algorithms" for massive parallelism and take account of all the constraints :-)

Right,

Are you trying to do a "raytracer" based on OpenCL ?

Thx

Great...

It is what I'm doing too...

Doing intersection per primitive is slow...so currently I'm working on implementing the acceleration structure with OpenCL... in order to handle millions of polygons ;-)

Take a look at :
- http://www.cs.unc.edu/~lauterb/GPUBVH/paper.pdf
- http://www.tml.tkk.fi/~timo/publications/aila2009hpg_paper.pdf

I haven't seen the example code yet, but this is very slow code:

Vector4 Normal = Vector4.Normalize( new Vector4( Vector3.Cross( EdgeAB.Xyz, EdgeAC.Xyz ), 0 ) );.

It constructs and copies 5 temporary vectors (the Xyz properties, the Cross and Normalize methods and the Vector4 constructor).

Rearranging the code to reduce the number of temporaries will have a positive impact to performance. Using OpenTK methods that take ref overloads will also help a lot:

float uw = Vector4.Dot(EdgeAB, w); // slow
 
float uw;
Vector4.Dot(ref EdgeAB, ref w, out uw); // fast

Unfortunately, it is very hard to write performant math code in .Net (this is a common complaint against XNA too). The lack of const references means you have to choose between multiple temporary objects and ugly syntax. (Even worse, sometimes it's impossible to avoid tempoararies.) Then, there's the lack of generic operators and other similar annoyances that make math code very tedious to implement.

I believe that even a CPU-powered OpenCL implementation will be faster for non-trivial applications.

I would say that CL vector operations are better optimized even on a CPU. There's no match for them in the managed world unless you get down to assembler/SSE. Still, the gap is so big that this is not the only difference. Depending on the implementation of OpenCL, multithreading capabilities of CPUs are exploited darn well, I'd say.

Right now, the only way of trying out the CPU vs GPU test is this:
1) First run the test on a CL capable nVidia card.
2) Uninstall nVidia drivers and run the test with ATI Stream on the CPU.

As of now, clGetPlatformIDs will not list both of these platforms. The vendors haven't reached an agreement in interoperability, yet. So we're stuck in one platform for now.

But you could still use multiple devices in one platform, right? Well, not so fast:
1) nVidia needs to cooperate with CPU vendors to release CPU capable OpenCL. [Edit: never mind the sentence here. I'm still asleep]
2) ATI releases GPU capable drivers. I love their HW since it's usually ahead of the competition. But the drivers have been a pain in the arse since I was a kid.
3) Get a MAC. It looks like you can run kernels on multiple devices with their integrated drivers. Not having one, I cannot confirm this.

2) AMD Stream 2.0 beta 4 is both CPU- and GPU-capable, so that's a viable route.

If you have a Nvidia GPU, there might be a way to hack a CPU/GPU testing environment together: first install AMD's Stream SDK (CPU only). Afterwards, install Nvidia's OpenCL-capable drivers (GPU only). If you run OpenCL now, you should get Nvidia's GPU implementation. If you wish to test with AMD's CPU implementation, simply copy the relevant dlls and exes from %Program Files%\Amd\Stream 2.0 to the test program's folder. To go back to Nvidia, simply rename AMD's OpenCL.dll to something else.

Edit: about that patent, I wouldn't worry about it too much. There's a *lot* of prior art (the oldest I could find is from 2004: http://www.codeproject.com/KB/cs/genericnumerics.aspx, but it's likely you can find even older implementations) and it probably won't hold up in court.

Extremely interesting performance comparison sample!

The AOBench kernel for OpenCL has the only output parameter of type global uint * result.

1. What is the size of this array for default picture size 256x256? 65536 I guess, isn't it?

2. How should I interpret this uint[] output to reproduce the picture? I am able to get it using AOBench for C#, and I wish to reproduce it for OpenCL implementation as well (just to be sure that both results are exactly identical).

After this step I will be ready to contribute AOBenchTest for Cloo (if anybody interested =) ).

Best regards,
Anton.

PS What is about Cloo 0.3.x announced a few days ago?