忍者ブログ

Memeplexes

プログラミング、3DCGとその他いろいろについて

[PR]

×

[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。

GPUでニューラルネットワーク更新（マトリクス×ベクトル）　その２　（OpenCL）

パフォーマンス改善

この間、GPUでニューラルネットワークを更新しましたが、その実装は適当でした。
遅いのです。
その実装を流用して制限付きボルツマンマシン(Restricted Boltzmann Machine : RBM)を実装した所、CPUより遅いというありさまでした。
今回はパフォーマンスを改善したやり方でニューラルネットワークの更新をしてみようと思います。
と言っても中身はただのマトリクスとベクトルの乗算です。
他にも流用できそうですね。


サンプルコード

ReductionNeuralNetwork.cs

using Cloo;
using System.Linq;

class ReductionNeuralNetwork
{
    const int visibleNeuronCount = 100;
    const int hiddenNeuronCount = 50 * 50;
    static float[] weights = Enumerable
        .Range(0, visibleNeuronCount * hiddenNeuronCount)
        .Select(i => (float)i)
        .ToArray();
    static float[] hiddenNeuronValues = Enumerable
        .Range(0, hiddenNeuronCount)
        .Select(i => (float)i)
        .ToArray();

    const int Multiply = 2;
    const int MinGroupSize = 256;

    static void Main()
    {
        var platform = ComputePlatform.Platforms[0];
        var devices = platform
            .Devices
            .Where(d => d.Type == ComputeDeviceTypes.Gpu)
            .ToArray();
        var context = new ComputeContext(
            devices,
            new ComputeContextPropertyList(platform),
            null,
            System.IntPtr.Zero
            );
        var commandQueue = new ComputeCommandQueue(
            context,
            devices[0],
            ComputeCommandQueueFlags.None
            );
        var program = new ComputeProgram(
            context,
            System.IO.File.ReadAllText("reduction.cl")
            );
        try
        {
            program.Build(devices, null, null, System.IntPtr.Zero);
        }
        catch
        {
            System.Console.WriteLine(program.GetBuildLog(devices[0]));
        }

        // updateVisibleNeurons
        var updateVisibleNeuronsKernel = program.CreateKernel("updateVisibleNeurons");
        var visibleNeuronInputStride = getNeuronInputStride(hiddenNeuronCount);
        var visibleNeuronBiasBuffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer,
            new float[visibleNeuronCount]
            );
        var weightBuffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer,
            weights
            );
        var hiddenNeuronValueBuffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer,
            Enumerable.Range(0, hiddenNeuronCount).Select(i => (float)i).ToArray()
            );
        var groupSize = System.Math.Min(
            MinGroupSize,
            updateVisibleNeuronsKernel.GetWorkGroupSize(devices[0])
            );
        var visibleNeuronValueBuffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite,
            visibleNeuronCount
            );
        var intermediateResultBuffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.CopyHostPointer,
            Enumerable.Range(0, (int)(visibleNeuronCount * (visibleNeuronInputStride / Multiply / groupSize))).Select(i => 0f).ToArray()
            );
        updateVisibleNeuronsKernel.SetMemoryArgument(0, visibleNeuronValueBuffer);
        updateVisibleNeuronsKernel.SetMemoryArgument(1, intermediateResultBuffer);
        updateVisibleNeuronsKernel.SetMemoryArgument(2, visibleNeuronBiasBuffer);
        updateVisibleNeuronsKernel.SetMemoryArgument(3, weightBuffer);
        updateVisibleNeuronsKernel.SetMemoryArgument(4, hiddenNeuronValueBuffer);
        updateVisibleNeuronsKernel.SetValueArgument(5, hiddenNeuronCount);
        updateVisibleNeuronsKernel.SetLocalArgument(6, groupSize * sizeof(float));

        var stopwatch = new System.Diagnostics.Stopwatch();
        stopwatch.Start();

        for (int i = 0; i < 100; i++)
        {
            commandQueue.Execute(
                updateVisibleNeuronsKernel,
                null,
                new long[] { visibleNeuronCount, visibleNeuronInputStride / Multiply },
                new long[] { 1, groupSize },
                null
                );
        }

        commandQueue.Finish();

        stopwatch.Stop();
        System.Console.WriteLine("calculation on gpu : " + stopwatch.Elapsed.TotalMilliseconds + "[ms]");
        /*System.Console.WriteLine("visible neuron inputs (cpu)");
        writeNeuronMatrix(getNeuronInputsByCpu(), hiddenNeuronCount);
        System.Console.WriteLine("intermediate result (gpu)");
        var intermediate = read(commandQueue, intermediateResultBuffer);

        for (int i = 0; i < visibleNeuronCount; i++)
        {
            for (int j = 0; j < intermediate.Length / visibleNeuronCount; j++)
            {
                System.Console.Write(
                    intermediate[i * intermediate.Length / visibleNeuronCount + j] + " "
                    );
            }

            System.Console.WriteLine();
        }

        System.Console.WriteLine("visible neuron values (gpu)");

        foreach (var number in read(commandQueue, visibleNeuronValueBuffer))
        {
            System.Console.Write(number + " ");
        }

        System.Console.WriteLine();
        System.Console.WriteLine("visible neuron values (cpu)");

        foreach (var number in getNeuronValuesByCpu())
        {
            System.Console.Write(number + " ");
        }*/
    }

    private static T[] read<T>(
        ComputeCommandQueue commandQueue,
        ComputeBuffer<T> buffer)
        where T : struct
    {
        var result = new T[buffer.Count];
        commandQueue.ReadFromBuffer(buffer, ref result, true, null);
        return result;
    }

    private static int getNeuronInputStride(int elementCount)
    {
        int MulFactor = Multiply * MinGroupSize;
        return (((elementCount - 1) / MulFactor) + 1) * MulFactor;
    }

    private static float[] getNeuronInputsByCpu()
    {
        var result = new float[visibleNeuronCount * hiddenNeuronCount];

        for (int i = 0; i < visibleNeuronCount; i++)
        {
            for (int j = 0; j < hiddenNeuronCount; j++)
            {
                result[i * hiddenNeuronCount + j] 
                    += weights[i * hiddenNeuronCount + j] 
                    * hiddenNeuronValues[j];
            }
        }

        return result;
    }

    private static void writeNeuronMatrix(float[] neuronMatrix, int stride)
    {
        for (int i = 0; i < visibleNeuronCount; i++)
        {
            for (int j = 0; j < hiddenNeuronCount; j++)
            {
                System.Console.Write(neuronMatrix[i * stride + j] + " ");
            }

            System.Console.WriteLine();
        }
    }

    private static float[] getNeuronValuesByCpu()
    {
        var result = new float[visibleNeuronCount];

        for (int i = 0; i < visibleNeuronCount; i++)
        {
            var sum = 0f;

            for (int j = 0; j < hiddenNeuronCount; j++)
            {
                sum  += weights[i * hiddenNeuronCount + j]
                    * hiddenNeuronValues[j];
            }

            result[i] = sum;
        }

        return result;
    }
}

はい、読みにくいです。
訳の分からないコメントアウトといいもうイヤになりますね。
これはParallel Reductionというアルゴリズムを利用しています。
おかげで読みにくいのですが、パフォーマンスはだいぶ改善されています。

reduction.cl

float sigmoid(float x)
{
	return 1.0f / (1.0f + exp(-x));
}

__kernel void updateVisibleNeurons(
	__global float *resultVisibleNeuronValues,
	__global float *intermediateResults,
	const __global float *visibleNeuronBiases, 
	const __global float *weights, 
	const __global float *hiddenNeuronValues,
	int hiddenNeuronCount,
	__local float *localData
	)
{
	int visibleNeuronIndex = get_global_id(0);
    int hiddenNeuronIndex = get_global_id(1);
    int localID = get_local_id(1);
	int visibleNeuronInputStride = get_global_size(1) * 2;
	int visibleNeuronOffset = visibleNeuronIndex * visibleNeuronInputStride;

	float inputs[2];
	inputs[0] = (hiddenNeuronIndex * 2 < hiddenNeuronCount)
		? 
		weights[visibleNeuronIndex * hiddenNeuronCount + hiddenNeuronIndex * 2] 
		* hiddenNeuronValues[hiddenNeuronIndex * 2]
		: 0;
	inputs[1] = (hiddenNeuronIndex * 2 + 1 < hiddenNeuronCount)
		?
		weights[visibleNeuronIndex * hiddenNeuronCount + hiddenNeuronIndex * 2 + 1] 
		* hiddenNeuronValues[hiddenNeuronIndex * 2 + 1]
		: 0;
	localData[localID] = inputs[0] + inputs[1];
    barrier(CLK_LOCAL_MEM_FENCE);

    for(int s = get_local_size(1) / 2; 0 < s; s /= 2) 
    {
        if(localID < s) 
        {
            localData[localID] += localData[localID + s];
        }

        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if(localID == 0)
	{
		intermediateResults[visibleNeuronIndex * get_num_groups(1) + get_group_id(1)] = localData[0];
	}

	barrier(CLK_GLOBAL_MEM_FENCE);

	if(hiddenNeuronIndex == 0)
	{
		float sum = 0;

		for(int i = 0; i < get_num_groups(1); i++)
		{
			sum += intermediateResults[visibleNeuronIndex * get_num_groups(1) + i];
		}

		resultVisibleNeuronValues[visibleNeuronIndex] = sum;
		//sigmoid(sum);
	}
}

シグモイド関数をコメントアウトで外しているので、これはただのマトリクスとベクトルのかけ算です。
実際にニューラルネットワークのシミュレーションをするときにはシグモイド関数を入れてください。

実行結果

これを実行すると、私の環境ではこうなりました：

calculation on gpu : 89.7232[ms]

だいたい0.1秒程度ですね。
ではこの間のアルゴリズムではどうでしょう？

calculation on gpu : 932.3391[ms]

約１秒です。
つまり、この間と比べておおよそ10倍速くなったのです！
よかったですね。

拍手[0回]

PR