-
Notifications
You must be signed in to change notification settings - Fork 32
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge pull request #65 from altimesh/november_update
November update
- Loading branch information
Showing
39 changed files
with
2,034 additions
and
91 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
9 changes: 9 additions & 0 deletions
9
HybridizerBasicSamples_CUDA100/7.Advanced/GenericReduction/GenericReduction/App.config
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,9 @@ | ||
| <?xml version="1.0" encoding="utf-8"?> | ||
| <configuration> | ||
| <startup> | ||
| <supportedRuntime version="v4.0" sku=".NETFramework,Version=v4.5.2" /> | ||
| </startup> | ||
| <appSettings> | ||
| <add key="hybridizer.cudaruntimeversion" value="100" /> | ||
| </appSettings> | ||
| </configuration> |
55 changes: 55 additions & 0 deletions
55
...BasicSamples_CUDA100/7.Advanced/GenericReduction/GenericReduction/GenericReduction.csproj
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,55 @@ | ||
| <?xml version="1.0" encoding="utf-8"?> | ||
| <Project ToolsVersion="14.0" DefaultTargets="Build" xmlns="http://schemas.microsoft.com/developer/msbuild/2003"> | ||
| <Import Project="$(MSBuildExtensionsPath)\$(MSBuildToolsVersion)\Microsoft.Common.props" Condition="Exists('$(MSBuildExtensionsPath)\$(MSBuildToolsVersion)\Microsoft.Common.props')" /> | ||
| <PropertyGroup> | ||
| <Configuration Condition=" '$(Configuration)' == '' ">Debug</Configuration> | ||
| <Platform Condition=" '$(Platform)' == '' ">x64</Platform> | ||
| <ProjectGuid>{4B7875CD-031E-40A4-ACF7-AE84124AA87F}</ProjectGuid> | ||
| <OutputType>Exe</OutputType> | ||
| <AppDesignerFolder>Properties</AppDesignerFolder> | ||
| <RootNamespace>GenericReduction</RootNamespace> | ||
| <AssemblyName>GenericReduction</AssemblyName> | ||
| <TargetFrameworkVersion>v4.0</TargetFrameworkVersion> | ||
| <FileAlignment>512</FileAlignment> | ||
| <AutoGenerateBindingRedirects>true</AutoGenerateBindingRedirects> | ||
| </PropertyGroup> | ||
| <PropertyGroup Condition="'$(Configuration)|$(Platform)' == 'Debug|x64'"> | ||
| <DebugSymbols>true</DebugSymbols> | ||
| <OutputPath>bin\x64\Debug\</OutputPath> | ||
| <DefineConstants>DEBUG;TRACE</DefineConstants> | ||
| <DebugType>full</DebugType> | ||
| <PlatformTarget>x64</PlatformTarget> | ||
| <ErrorReport>prompt</ErrorReport> | ||
| <CodeAnalysisRuleSet>MinimumRecommendedRules.ruleset</CodeAnalysisRuleSet> | ||
| <Prefer32Bit>true</Prefer32Bit> | ||
| </PropertyGroup> | ||
| <PropertyGroup Condition="'$(Configuration)|$(Platform)' == 'Release|x64'"> | ||
| <OutputPath>bin\x64\Release\</OutputPath> | ||
| <DefineConstants>TRACE</DefineConstants> | ||
| <Optimize>true</Optimize> | ||
| <DebugType>pdbonly</DebugType> | ||
| <PlatformTarget>x64</PlatformTarget> | ||
| <ErrorReport>prompt</ErrorReport> | ||
| <CodeAnalysisRuleSet>MinimumRecommendedRules.ruleset</CodeAnalysisRuleSet> | ||
| <Prefer32Bit>true</Prefer32Bit> | ||
| </PropertyGroup> | ||
| <ItemGroup> | ||
| <Reference Include="System" /> | ||
| <Reference Include="System.Core" /> | ||
| <Reference Include="System.Xml.Linq" /> | ||
| <Reference Include="System.Data.DataSetExtensions" /> | ||
| <Reference Include="Microsoft.CSharp" /> | ||
| <Reference Include="System.Data" /> | ||
| <Reference Include="System.Net.Http" /> | ||
| <Reference Include="System.Xml" /> | ||
| <Reference Include="Hybridizer.Runtime.CUDAImports" /> | ||
| </ItemGroup> | ||
| <ItemGroup> | ||
| <Compile Include="Program.cs" /> | ||
| <Compile Include="Properties\AssemblyInfo.cs" /> | ||
| </ItemGroup> | ||
| <ItemGroup> | ||
| <None Include="App.config" /> | ||
| </ItemGroup> | ||
| <Import Project="$(MSBuildToolsPath)\Microsoft.CSharp.targets" /> | ||
| </Project> |
163 changes: 163 additions & 0 deletions
163
HybridizerBasicSamples_CUDA100/7.Advanced/GenericReduction/GenericReduction/Program.cs
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,163 @@ | ||
| using Hybridizer.Runtime.CUDAImports; | ||
| using System; | ||
| using System.Collections.Generic; | ||
| using System.Linq; | ||
| using System.Text; | ||
| using System.Threading; | ||
| using System.Threading.Tasks; | ||
|
|
||
| namespace GenericReduction | ||
| { | ||
| [HybridTemplateConcept] | ||
| interface IReductor | ||
| { | ||
| [Kernel] | ||
| float func(float x, float y); | ||
| [Kernel] | ||
| float neutral { get; } | ||
| } | ||
|
|
||
| struct AddReductor: IReductor | ||
| { | ||
| [Kernel] | ||
| public float neutral { get { return 0.0F; } } | ||
|
|
||
| [Kernel] | ||
| public float func(float x, float y) | ||
| { | ||
| return x + y; | ||
| } | ||
| } | ||
|
|
||
| struct MaxReductor : IReductor | ||
| { | ||
| [Kernel] | ||
| public float neutral { get { return float.MinValue; } } | ||
|
|
||
| [Kernel] | ||
| public float func(float x, float y) | ||
| { | ||
| return Math.Max(x, y); | ||
| } | ||
| } | ||
|
|
||
| [HybridRegisterTemplate(Specialize = typeof(GridReductor<MaxReductor>))] | ||
| [HybridRegisterTemplate(Specialize = typeof(GridReductor<AddReductor>))] | ||
| class GridReductor<TReductor> where TReductor : struct, IReductor | ||
| { | ||
| [Kernel] | ||
| TReductor reductor { get { return default(TReductor); } } | ||
|
|
||
| [Kernel] | ||
| public void Reduce(float[] result, float[] input, int N) | ||
| { | ||
| var cache = new SharedMemoryAllocator<float>().allocate(blockDim.x); | ||
| int tid = threadIdx.x + blockDim.x * blockIdx.x; | ||
| int cacheIndex = threadIdx.x; | ||
|
|
||
| float tmp = reductor.neutral; | ||
| while (tid < N) | ||
| { | ||
| tmp = reductor.func(tmp, input[tid]); | ||
| tid += blockDim.x * gridDim.x; | ||
| } | ||
|
|
||
| cache[cacheIndex] = tmp; | ||
|
|
||
| CUDAIntrinsics.__syncthreads(); | ||
|
|
||
| int i = blockDim.x / 2; | ||
| while (i != 0) | ||
| { | ||
| if (cacheIndex < i) | ||
| { | ||
| cache[cacheIndex] = reductor.func(cache[cacheIndex], cache[cacheIndex + i]); | ||
| } | ||
|
|
||
| CUDAIntrinsics.__syncthreads(); | ||
| i >>= 1; | ||
| } | ||
|
|
||
| if (cacheIndex == 0) | ||
| { | ||
| result[blockIdx.x] = cache[0]; | ||
| } | ||
| } | ||
| } | ||
|
|
||
| // Unfortunately this is necessay since we didn't implemented generic entrypoints yet. | ||
| class EntryPoints | ||
| { | ||
| [EntryPoint] | ||
| public static void ReduceAdd(GridReductor<AddReductor> reductor, float[] result, float[] input, int N) | ||
| { | ||
| reductor.Reduce(result, input, N); | ||
| } | ||
| [EntryPoint] | ||
| public static void ReduceMax(GridReductor<MaxReductor> reductor, float[] result, float[] input, int N) | ||
| { | ||
| reductor.Reduce(result, input, N); | ||
| } | ||
| } | ||
|
|
||
| class Program | ||
| { | ||
| static void Main(string[] args) | ||
| { | ||
| const int N = 1024 * 1024 * 32; | ||
| float[] a = new float[N]; | ||
|
|
||
| // initialization | ||
| float cudaMax = float.MinValue; | ||
| float cudaAdd = 0.0F; | ||
| Random random = new Random(42); | ||
| Parallel.For(0, N, i => a[i] = (float)random.NextDouble()); | ||
|
|
||
| // hybridizer configuration | ||
| cudaDeviceProp prop; | ||
| cuda.GetDeviceProperties(out prop, 0); | ||
| int gridDimX = 16 * prop.multiProcessorCount; | ||
| HybRunner runner = HybRunner.Cuda().SetDistrib(gridDimX, 1, 256, 1, 1, gridDimX * sizeof(float)); | ||
| float[] buffMax = new float[gridDimX]; | ||
| float[] buffAdd = new float[gridDimX]; | ||
| var maxReductor = new GridReductor<MaxReductor>(); | ||
| var addReductor = new GridReductor<AddReductor>(); | ||
| dynamic wrapped = runner.Wrap(new EntryPoints()); | ||
|
|
||
| // device reduction | ||
| wrapped.ReduceMax(maxReductor, buffMax, a, N); | ||
| wrapped.ReduceAdd(addReductor, buffAdd, a, N); | ||
| cuda.ERROR_CHECK(cuda.DeviceSynchronize()); | ||
|
|
||
| // finalize reduction host side | ||
| for (int i = 0; i < gridDimX; ++i) | ||
| { | ||
| cudaMax = Math.Max(cudaMax, buffMax[i]); | ||
| cudaAdd += buffAdd[i]; | ||
| } | ||
|
|
||
| // check results | ||
| float expectedMax = a.AsParallel().Aggregate((x, y) => Math.Max(x, y)); | ||
| float expectedAdd = a.AsParallel().Aggregate((x, y) => x + y); | ||
| bool hasError = false; | ||
| if (cudaMax != expectedMax) | ||
| { | ||
| Console.Error.WriteLine($"MAX Error : {cudaMax} != {expectedMax}"); | ||
| hasError = true; | ||
| } | ||
|
|
||
| // addition is not associative, so results cannot be exactly the same | ||
| // https://en.wikipedia.org/wiki/Associative_property#Nonassociativity_of_floating_point_calculation | ||
| if (Math.Abs(cudaAdd - expectedAdd) / expectedAdd > 1.0E-5F) | ||
| { | ||
| Console.Error.WriteLine($"ADD Error : {cudaAdd} != {expectedAdd}"); | ||
| hasError = true; | ||
| } | ||
|
|
||
| if (hasError) | ||
| Environment.Exit(1); | ||
|
|
||
| Console.Out.WriteLine("OK"); | ||
| } | ||
| } | ||
| } |
41 changes: 41 additions & 0 deletions
41
...icSamples_CUDA100/7.Advanced/GenericReduction/GenericReduction/Properties/AssemblyInfo.cs
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,41 @@ | ||
| using System.Reflection; | ||
| using System.Runtime.CompilerServices; | ||
| using System.Runtime.InteropServices; | ||
|
|
||
| using Hybridizer.Runtime.CUDAImports; | ||
|
|
||
| // General Information about an assembly is controlled through the following | ||
| // set of attributes. Change these attribute values to modify the information | ||
| // associated with an assembly. | ||
| [assembly: AssemblyTitle("GenericReduction")] | ||
| [assembly: AssemblyDescription("")] | ||
| [assembly: AssemblyConfiguration("")] | ||
| [assembly: AssemblyCompany("")] | ||
| [assembly: AssemblyProduct("GenericReduction")] | ||
| [assembly: AssemblyCopyright("Copyright © 2018")] | ||
| [assembly: AssemblyTrademark("")] | ||
| [assembly: AssemblyCulture("")] | ||
|
|
||
| // Setting ComVisible to false makes the types in this assembly not visible | ||
| // to COM components. If you need to access a type in this assembly from | ||
| // COM, set the ComVisible attribute to true on that type. | ||
| [assembly: ComVisible(false)] | ||
|
|
||
| // The following GUID is for the ID of the typelib if this project is exposed to COM | ||
| [assembly: Guid("4b7875cd-031e-40a4-acf7-ae84124aa87f")] | ||
|
|
||
| // Version information for an assembly consists of the following four values: | ||
| // | ||
| // Major Version | ||
| // Minor Version | ||
| // Build Number | ||
| // Revision | ||
| // | ||
| // You can specify all the values or you can default the Build and Revision Numbers | ||
| // by using the '*' as shown below: | ||
| // [assembly: AssemblyVersion("1.0.*")] | ||
| [assembly: AssemblyVersion("1.0.0.0")] | ||
| [assembly: AssemblyFileVersion("1.0.0.0")] | ||
|
|
||
| // allow parameterless constructor for HybRunner | ||
| [assembly: HybRunnerDefaultSatelliteName("GenericReduction_CUDA.dll")] |
Oops, something went wrong.