Skip to content

Commit 88926c0

Browse files
deepakkumar1984deepakkumar1984
authored
Merge pull request #8 from kgamecarter/master
Add half type, vload vstore function
2 parents dafa842 + 274c0cc commit 88926c0

File tree

4 files changed

+650
-4
lines changed

4 files changed

+650
-4
lines changed

‎examples/AmplifierExamples/Kernels/SimpleKernels.cs‎

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,14 +8,23 @@ namespace AmplifierExamples.Kernels
88
class SimpleKernels : OpenCLFunctions
99
{
1010
[OpenCLKernel]
11-
void AddData([Global, Input]float[] a, [Global] float[] b, [Global, Output]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];
1616
a[i] += 2; // result will not copy out
1717
}
1818

19+
[OpenCLKernel]
20+
void AddHalf([Global, Input]half[] a, [Global]half[] b)
21+
{
22+
int i = get_global_id(0);
23+
float af = vload_half(i, a);
24+
float bf = vload_half(i, b);
25+
vstore_half(af + bf, i, b);
26+
}
27+
1928
[OpenCLKernel]
2029
void Fill([Global] float[] x, float value)
2130
{

‎examples/AmplifierExamples/SimpleKernelEx.cs‎

Lines changed: 10 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
11
using Amplifier;
2+
using Amplifier.OpenCL;
23
using AmplifierExamples.Kernels;
34
using System;
45
using System.Collections.Generic;
@@ -34,9 +35,9 @@ public void Execute()
3435
}
3536

3637
//Create variable a, b and r
37-
Array x = new float[] { 1, 2, 3, 4, 5, 6, 7, 8, 9 };
38-
Array y = new float[9];
39-
Array z = new float[9];
38+
var x = new float[] { 1, 2, 3, 4, 5, 6, 7, 8, 9 };
39+
var y = new float[9];
40+
var z = new float[9];
4041

4142
//Get the execution engine
4243
var exec = compiler.GetExec();
@@ -47,6 +48,12 @@ public void Execute()
4748
//Execute AddData kernel method
4849
exec.AddData(x, y, z);
4950

51+
//Execute AddHalf kernel method
52+
var xhalf = Array.ConvertAll(x, v => (half)v);
53+
var yhalf = Array.ConvertAll(y, v => (half)v);
54+
exec.AddHalf(xhalf, yhalf);
55+
z = Array.ConvertAll(yhalf, v => (float)v);
56+
5057
//Execuete SAXPY kernel method
5158
exec.SAXPY(x, y, 2f);
5259

‎src/Amplifier.Net/OpenCL/DataTypes/ScalarDataTypes.cs‎

Lines changed: 122 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,4 +35,126 @@ public static implicit operator uintptr_t(uint d)
3535
return new uintptr_t();
3636
}
3737
}
38+
39+
public struct half
40+
{
41+
private ushort Value;
42+
43+
public override string ToString()
44+
{
45+
return ((float)this).ToString();
46+
}
47+
48+
public static explicit operator half(float d)
49+
{
50+
return new half(d);
51+
}
52+
53+
public unsafe static explicit operator float(half d)
54+
{
55+
bool isPos = (d.Value & Float16Params.SignMask) == 0;
56+
uint biasedExponent = (d.Value & Float16Params.ExpMask) >> Float16Params.ExpOffset;
57+
uint frac = (d.Value & Float16Params.FracMask);
58+
bool isInf = biasedExponent == Float16Params.BiasedExpMax && (frac == 0);
59+
60+
if (isInf)
61+
{
62+
return isPos ? float.PositiveInfinity : float.NegativeInfinity;
63+
}
64+
65+
bool isNan = biasedExponent == Float16Params.BiasedExpMax && (frac != 0);
66+
if (isNan)
67+
{
68+
return float.NaN;
69+
}
70+
71+
bool isSubnormal = biasedExponent == 0;
72+
if (isSubnormal)
73+
{
74+
return frac * Float16Params.SmallestSubnormalAsFloat * (isPos ? 1.0f : -1.0f);
75+
}
76+
77+
int unbiasedExp = (int)biasedExponent - Float16Params.ExpBias;
78+
uint biasedF32Exponent = (uint)(unbiasedExp + Float32Params.ExpBias);
79+
80+
uint bits;
81+
82+
bits = (isPos ? 0u : 1u << Float32Params.SignOffset)
83+
| (biasedF32Exponent << Float32Params.ExpOffset)
84+
| (frac << (Float32Params.ExpOffset - Float16Params.ExpOffset));
85+
86+
return *(float*)&bits;
87+
}
88+
89+
public unsafe half(float d)
90+
{
91+
uint bits = *(uint*)&d;
92+
93+
uint fAbsBits = bits & Float32Params.AbsValueMask;
94+
bool isNeg = (bits & Float32Params.SignBitMask) != 0;
95+
uint sign = (bits & Float32Params.SignBitMask) >> (Float16Params.NumFracBits + Float16Params.NumExpBits + 1);
96+
uint half;
97+
98+
if (float.IsNaN(d))
99+
{
100+
half = (Float16Params.ExpMask | Float16Params.FracMask);
101+
}
102+
else if (float.IsInfinity(d))
103+
{
104+
half = isNeg ? Float16Params.SignMask | Float16Params.ExpMask : Float16Params.ExpMask;
105+
}
106+
else if (fAbsBits > Float16Params.MaxNormal)
107+
{
108+
// Clamp to max float 16 value
109+
half = sign | (((1 << Float16Params.NumExpBits) - 1) << Float16Params.NumFracBits) | Float16Params.FracMask;
110+
}
111+
else if (fAbsBits < Float16Params.MinNormal)
112+
{
113+
uint fracBits = (fAbsBits & Float32Params.MantissaMask) | (1 << Float32Params.NumMantissaBits);
114+
int nshift = Float16Params.Emin + Float32Params.Emax - (int)(fAbsBits >> Float32Params.NumMantissaBits);
115+
uint shiftedBits = nshift < 24 ? fracBits >> nshift : 0;
116+
half = sign | (shiftedBits >> Float16Params.FracBitsDiff);
117+
}
118+
else
119+
{
120+
half = sign | ((fAbsBits + Float16Params.BiasDiff) >> Float16Params.FracBitsDiff);
121+
}
122+
this.Value = (ushort)half;
123+
}
124+
125+
private static class Float16Params
126+
{
127+
public const uint BitSize = 16; // total number of bits in the representation
128+
public const int NumFracBits = 10; // number of fractional (mantissa) bits
129+
public const int NumExpBits = 5; // number of (biased) exponent bits
130+
public const uint SignBit = 15; // position of the sign bit
131+
public const uint SignMask = 1 << 15; // mask to extract sign bit
132+
public const uint FracMask = (1 << 10) - 1; // mask to extract the fractional (mantissa) bits
133+
public const uint ExpMask = ((1 << 5) - 1) << 10; // mask to extract the exponent bits
134+
public const uint Emax = (1 << (5 - 1)) - 1; // max value for the exponent
135+
public const int Emin = -((1 << (5 - 1)) - 1) + 1; // min value for the exponent
136+
public const uint MaxNormal = ((((1 << (5 - 1)) - 1) + 127) << 23) | 0x7FE000; // max value that can be represented by the 16 bit float
137+
public const uint MinNormal = ((-((1 << (5 - 1)) - 1) + 1) + 127) << 23; // min value that can be represented by the 16 bit float
138+
public const uint BiasDiff = unchecked((uint)(((1 << (5 - 1)) - 1) - 127) << 23); // difference in bias between the float16 and float32 exponent
139+
public const int FracBitsDiff = 23 - 10; // difference in number of fractional bits between float16/float32
140+
141+
public const int ExpBias = 15;
142+
public const int ExpOffset = 10;
143+
public const ushort BiasedExpMax = (1 << 5) - 1;
144+
public const float SmallestSubnormalAsFloat = 5.96046448e-8f;
145+
}
146+
147+
private static class Float32Params
148+
{
149+
public const uint AbsValueMask = 0x7FFFFFFF; // ANDing with this value gives the abs value
150+
public const uint SignBitMask = 0x80000000; // ANDing with this value gives the sign
151+
public const int Emax = 127; // max value for the exponent
152+
public const int NumMantissaBits = 23; // 23 bit mantissa on single precision floats
153+
public const uint MantissaMask = 0x007FFFFF; // 23 bit mantissa on single precision floats
154+
155+
public const int SignOffset = 31;
156+
public const int ExpBias = 127;
157+
public const int ExpOffset = 23;
158+
}
159+
}
38160
}

0 commit comments

Comments
 (0)