Skip to content

Commit dafa842

Browse files
deepakkumar1984deepakkumar1984
authored
Merge pull request #7 from kgamecarter/master
Add parameter of kernel function is only input or output
2 parents 954ec38 + afcdb75 commit dafa842

File tree

7 files changed

+88
-16
lines changed

7 files changed

+88
-16
lines changed

‎examples/AmplifierExamples/Kernels/SimpleKernels.cs‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,11 +8,12 @@ namespace AmplifierExamples.Kernels
88
class SimpleKernels : OpenCLFunctions
99
{
1010
[OpenCLKernel]
11-
void AddData([Global]float[] a, [Global] float[] b, [Global]float[] r)
11+
void AddData([Global, Input]float[] a, [Global] float[] b, [Global, Output]float[] r)
1212
{
1313
int i = get_global_id(0);
1414
b[i] = 0.5f * b[i];
1515
r[i] = a[i] + b[i];
16+
a[i] += 2; // result will not copy out
1617
}
1718

1819
[OpenCLKernel]

‎examples/AmplifierExamples/SimpleKernelEx.cs‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -36,13 +36,17 @@ public void Execute()
3636
//Create variable a, b and r
3737
Array x = new float[] { 1, 2, 3, 4, 5, 6, 7, 8, 9 };
3838
Array y = new float[9];
39+
Array z = new float[9];
3940

4041
//Get the execution engine
4142
var exec = compiler.GetExec();
4243

4344
//Execute fill kernel method
4445
exec.Fill(y, 0.5f);
4546

47+
//Execute AddData kernel method
48+
exec.AddData(x, y, z);
49+
4650
//Execuete SAXPY kernel method
4751
exec.SAXPY(x, y, 2f);
4852

‎src/Amplifier.Net/KernelFunction.cs‎

Lines changed: 28 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@ public class KernelFunction
1616
/// </summary>
1717
public KernelFunction()
1818
{
19-
Parameters = new Dictionary<string, string>();
19+
Parameters = new Dictionary<string, FunctionParameter>();
2020
}
2121

2222
/// <summary>
@@ -33,6 +33,32 @@ public KernelFunction()
3333
/// <value>
3434
/// The parameters.
3535
/// </value>
36-
public Dictionary<string, string> Parameters { get; set; }
36+
public Dictionary<string, FunctionParameter> Parameters { get; set; }
37+
}
38+
39+
/// <summary>
40+
/// Class to hold the parameter of kernel method info
41+
/// </summary>
42+
public class FunctionParameter
43+
{
44+
/// <summary>
45+
/// Gets or sets the type name.
46+
/// </summary>
47+
public string TypeName { get; set; }
48+
49+
/// <summary>
50+
/// Gets or sets the IO mode.
51+
/// </summary>
52+
public IOMode IOMode { get; set; }
53+
}
54+
55+
/// <summary>
56+
/// The parameter of kernel method is input or output
57+
/// </summary>
58+
public enum IOMode
59+
{
60+
InOut = 0,
61+
In,
62+
Out
3763
}
3864
}

‎src/Amplifier.Net/OpenCL/Cloo/GenericArrayMemory.cs‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,9 +23,14 @@ public GenericArrayMemory(ComputeContext context, ComputeMemoryFlags flags, obje
2323
return;
2424

2525
int size = Marshal.SizeOf(array.GetValue(0).GetType()) * array.Length;
26-
var datagch = GCHandle.Alloc(obj, GCHandleType.Pinned);
26+
var hostPtr = IntPtr.Zero;
27+
if ((flags & (ComputeMemoryFlags.CopyHostPointer | ComputeMemoryFlags.UseHostPointer)) != ComputeMemoryFlags.None)
28+
{
29+
var datagch = GCHandle.Alloc(obj, GCHandleType.Pinned);
30+
hostPtr = datagch.AddrOfPinnedObject();
31+
}
2732
ComputeErrorCode error = ComputeErrorCode.Success;
28-
var handle = CL12.CreateBuffer(context.Handle, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer, new IntPtr(size), datagch.AddrOfPinnedObject(), out error);
33+
var handle = CL12.CreateBuffer(context.Handle, flags, new IntPtr(size), hostPtr, out error);
2934

3035
this.Size = size;
3136
this.Handle = handle;

‎src/Amplifier.Net/OpenCL/CodeTranslator.cs‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,8 @@ private static string GlobalCleanup(string code)
3030
.Replace("[OpenCLKernel]", "__kernel")
3131
.Replace("[Global]", "global")
3232
.Replace("[Struct]", "struct")
33+
.Replace("[Input]", "")
34+
.Replace("[Output]", "")
3335
.Replace("@", "v_");
3436

3537
foreach (var item in replaceEmptyList)

‎src/Amplifier.Net/OpenCL/OpenCLAttribute.cs‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,16 @@ public class GlobalAttribute : System.Attribute
1010
{
1111
}
1212

13+
[System.AttributeUsage(System.AttributeTargets.Parameter)]
14+
public class InputAttribute : System.Attribute
15+
{
16+
}
17+
18+
[System.AttributeUsage(System.AttributeTargets.Parameter)]
19+
public class OutputAttribute : System.Attribute
20+
{
21+
}
22+
1323
[System.AttributeUsage(System.AttributeTargets.Parameter)]
1424
public class StructAttribute : System.Attribute
1525
{

‎src/Amplifier.Net/OpenCLCompiler.cs‎

Lines changed: 35 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -220,16 +220,22 @@ public override void Execute(string functionName, params object[] args)
220220
var ndobject = (Array)args.FirstOrDefault(x => (x.GetType().IsArray));
221221
long length = ndobject != null ? ndobject.Length : 1;
222222

223-
var buffers = BuildKernelArguments(args, kernel, length);
223+
var method = KernelFunctions.FirstOrDefault(x => (x.Name == functionName));
224+
225+
var buffers = BuildKernelArguments(method, args, kernel, length);
224226
commands.Execute(kernel, null, new long[] { length }, null, null);
225227

226-
for(int i=0;i<args.Length;i++)
228+
for (int i = 0; i < args.Length; i++)
227229
{
228230
if (!args[i].GetType().IsArray)
229231
continue;
230232

231-
Array r = (Array)args[i];
232-
commands.ReadFromMemory(buffers[i], ref r, true, 0, null);
233+
var ioMode = method.Parameters.ElementAt(i).Value.IOMode;
234+
if (ioMode == IOMode.InOut || ioMode == IOMode.Out)
235+
{
236+
Array r = (Array)args[i];
237+
commands.ReadFromMemory(buffers[i], ref r, true, 0, null);
238+
}
233239
buffers[i].Dispose();
234240
}
235241
}
@@ -277,11 +283,11 @@ private void ValidateArgs(string functionName, object[] args)
277283

278284
if (args[i].GetType().IsPrimitive)
279285
{
280-
args[i] = Convert.ChangeType(args[i], Type.GetType(parameter.Value));
286+
args[i] = Convert.ChangeType(args[i], Type.GetType(parameter.Value.TypeName));
281287
}
282288
else if (args[i].GetType().IsArray)
283289
{
284-
if (parameter.Value != args[i].GetType().FullName)
290+
if (parameter.Value.TypeName != args[i].GetType().FullName)
285291
throw new ExecutionException(string.Format("Data type mismatch for parameter {0}. Expected is {1} but got {2}",
286292
parameter.Key,
287293
(parameter.Value,
@@ -417,7 +423,20 @@ CSharpDecompiler cSharpDecompiler
417423
var k = new KernelFunction() { Name = item.Name };
418424
foreach (var p in item.Parameters)
419425
{
420-
k.Parameters.Add(p.Name, p.Type.FullName);
426+
var isInput = p.GetAttributes().Any(x => x.AttributeType.Name == "InputAttribute");
427+
var isOutput = p.GetAttributes().Any(x => x.AttributeType.Name == "OutputAttribute");
428+
var mode = IOMode.InOut;
429+
if (isInput)
430+
mode = IOMode.In;
431+
if (isOutput)
432+
mode = IOMode.Out;
433+
if (isInput && isOutput)
434+
mode = IOMode.InOut;
435+
k.Parameters.Add(p.Name, new FunctionParameter
436+
{
437+
TypeName = p.Type.FullName,
438+
IOMode = mode
439+
});
421440
}
422441

423442
KernelFunctions.Add(k);
@@ -469,12 +488,13 @@ private void CreateKernels(string code)
469488
/// Builds the kernel arguments.
470489
/// </summary>
471490
/// <typeparam name="TSource">The type of the source.</typeparam>
491+
/// <param name="method">The method.</param>
472492
/// <param name="inputs">The inputs.</param>
473493
/// <param name="kernel">The kernel.</param>
474494
/// <param name="length">The length.</param>
475495
/// <param name="returnInputVariable">The return result.</param>
476496
/// <returns></returns>
477-
private Dictionary<int, GenericArrayMemory> BuildKernelArguments(object[] inputs, ComputeKernel kernel, long length, int? returnInputVariable = null)
497+
private Dictionary<int, GenericArrayMemory> BuildKernelArguments(KernelFunction method, object[] inputs, ComputeKernel kernel, long length, int? returnInputVariable = null)
478498
{
479499
int i = 0;
480500
Dictionary<int, GenericArrayMemory> result = new Dictionary<int, GenericArrayMemory>();
@@ -484,9 +504,13 @@ private Dictionary<int, GenericArrayMemory> BuildKernelArguments(object[] inputs
484504
int size = 0;
485505
if(item.GetType().IsArray)
486506
{
487-
488-
var datagch = GCHandle.Alloc(item, GCHandleType.Pinned);
489-
GenericArrayMemory mem = new GenericArrayMemory(_context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer, item);
507+
var mode = method.Parameters.ElementAt(i).Value.IOMode;
508+
var flag = ComputeMemoryFlags.ReadWrite;
509+
if (mode == IOMode.Out)
510+
flag |= ComputeMemoryFlags.AllocateHostPointer;
511+
else
512+
flag |= ComputeMemoryFlags.CopyHostPointer;
513+
GenericArrayMemory mem = new GenericArrayMemory(_context, flag, item);
490514
kernel.SetMemoryArgument(i, mem);
491515
result.Add(i, mem);
492516
}

0 commit comments

Comments
 (0)