Memeplexes

[PR]

2025年02月26日

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

v4.0でv2.0のアセンブリを使うとFileLoadExceptionがスローされる問題 (Mixed mode assembly is built against version 'v2.0.50727' of the runtime and cannot be loaded in the 4.0 runtime without additional configuration information.)

このブログでも何度か取り上げましたが、独立して記事にしておきます。
VS2010などで、ランタイムバージョンが2.0のアセンブリを参照して使うと、次のような例外が出ます。

FileLoadExceptionです。
ようは、アセンブリのバージョンが違うのでだめですよと言っているのです。

この問題を解決するには、.configファイルを使います。
バージョン２も使いますよと教えてあげるのです。

App.config


<?xml version="1.0" encoding="utf-8" ?>
<configuration>
  <startup useLegacyV2RuntimeActivationPolicy="true">
    <supportedRuntime version="v4.0"/>
  </startup>
</configuration>

このファイルを作って、プログファムを実行するとさっきのような例外はスローされなくなるはずです。

SlimDXでは

さて、私がこの問題に直面したのはSlimDXを使っている時でした。
SlimDXとはDirectXのC#用ラッパーライブラリーです。

このSlimDXのアセンブリを参照に追加しようとしている時です。
どうも一見SlimDXにはランタイムバージョン4.0のものが存在しないように見えたのです。
あたかも2.0のアセンブリだけであるかのように。

クリックして上の画像を拡大してみてください。
SlimDXにはランタイムバージョンが2.0しか無いように、一見見えます。
ですから私はそのままこのアセンブリをv4.0のプロジェクトに追加して、FileLoadExceptionに悩まされました。
冒頭のようにApp.configを追加しなければならなかったのです。

そんなわけで「なんでSlimDXはv4.0対応していないのだろう？不親切だなー。これさえなければいいライブラリなのに・・・」と思ったものです。
が、違うのです。
v4.0のものもちゃんと有ります。

思い出すだけで恥ずかしいのですが、これはRuntimeというところをクリックすると、v4.0のアセンブリが出てくるのです。
v4.0のものが先頭に出てきてからSlimDXで検索すると、きちんと出てきます。

ですから、実はSlimDXを使うときにApp.configファイルを作る必要など全くなかったのです。
うーん反省です。

[2回]

Windows7 vs Windows XP（アクセス数~普及率）

ぼーっとアクセスを眺めていて気がついたのですが、最近このブログへのユニークアクセス数で、Windows7がXPを抜き、１位になりました。
他のところではどうだかわかりませんがともかく、
祝！世代交代！！

おわかりいただけたでしょうか。
このグラフはこのブログへアクセスしたPCをOS別に分類し、その推移を描いています。
アクセス数は普及率を反映していると思われ、このグラフはそのままOSの普及率として読み替えてもいいでしょう。

このとおり、３月まではXPが一位ですが、４月からここ３ヶ月はずっと7が一位です！
つまりXPの時代は終わり、7の時代が幕を開けたのです！！！

……
いや……
どうでしょうね。
記録がないのでわかりませんが、２月がどうだったのか。
２月がもし7が一位だったとすれば、世代交代はもうずっと前に起こっていた、あるいは今はまだ起こっている最中だと考えることもできそうです。

しかしまあ数だけ見れば今Windows 7が一位だということは確かです！
DirectXの対応状況を考えれば、これはとても好ましいこと。
現在の最新バージョンはDirectX11だと言うのに、XPはDirectX9までにしか対応していませんからね。
7が一位を取ったということは、それだけDirectX10以降の対応PCの数が増えている。
つまり、グラフィックスの進歩と前進につながるはずです・・・！！！！

ただ、７が一位を取ったというのには、おそらくこのブログでDirectX10や11を扱ったというのも、関係があるかもしれません。
つまりデータにバイアスがかかっているのです。
DirectX10以降を実行できないPCでは、このブログに書かれていることは、あまり意味が無いかもしれないですからね。
うーん・・・

[1回]

[CUDA] メモリーコアレッシング(Memory coalescing)

CUDAの話です。
先日CUDAのコードを読んでいると、不可解なコードに行き当たりました。

energygrid[outaddr] += energyvalx1;
energygrid[outaddr + 1 * BLOCKSIZEX] += energyvalx2;
energygrid[outaddr + 2 * BLOCKSIZEX] += energyvalx3;
...
energygrid[outaddr + 6 * BLOCKSIZEX] += energyvalx7;
energygrid[outaddr + 7 * BLOCKSIZEX] += energyvalx8;

これは配列energygrid[]に、floatの値を足しているようです。
それはいいのですが、インデックスに注目してください。
1 * BLOCKSIZEXとはなんでしょう？
2 * BLOCKSIZEXとは？

調べてみると、BLOCKSIZEXは16でした。
つまり言い換えると次のようにアクセスしていることになります。

energygrid[outaddr] += energyvalx1;
energygrid[outaddr + 16] += energyvalx2;
energygrid[outaddr + 32] += energyvalx3;
...
energygrid[outaddr + 96] += energyvalx7;
energygrid[outaddr + 112] += energyvalx8;

……
ちょっと待ってください。
なぜ１６飛び飛びにアクセスしているのでしょうか？
このコードの元の目的から考えると（ここでは解説しませんが）、プログラムはenergygridすべての要素にアクセスします。
というか全要素にアクセスしなければなりません。
全要素にアクセスするなら飛び飛びアクセスでなくてもいいんじゃ！？
単に順番に、[outaddr]、[outaddr + 1]、[outaddr + 2]、...[outaddr + 7]というふうにアクセスしてはなぜいけないのでしょう？

しばらく頭を捻った後ようやくわかりました。
つまりこれは複数のスレッドで考えると話が見えてくるのです。
今までは一つのスレッドでこう考えていました：

こう考えるとたしかに不可解なのです。
インデックスが１６飛び飛びですから。
「どうして0,1,2,3じゃなくて0, 16, 32, 48なの？」となるわけです。

しかしCUDAはGPGPU。
スレッドは多数用意して計算するもの。
そうすると実行時には次のようになります：

おわかりいただけたでしょうか…
一つのスレッドで考えるとたしかに飛び飛びだったものが、複数スレッドで見るとアクセスが綺麗に連続していることがわかります。
全部のスレッドの最初のアクセスは、0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15と連続しています。
そのつぎもやはり連続しています。

CUDAではこのように複数のスレッドの固まりが連続したところにアクセスしていると、そのメモリアクセスをまとめ（Coalescing）てくれるそうなのです。
こうしてまとめられたメモリアクセスは、そうでないものより効率がいいそうです。
だからわざわざ*16していたのですね。

ちなみに16というのは、半ワープというもののサイズです。
CUDAでは実行するスレッドはワープという固まりをなします。
それが32個です。
１ワープ３２スレッド。
半ワープというのはその半分の16個です。
メモリアクセスのまとめ（Coalescing）は、半ワープが単位となっているそうです。

なお、GPU使っているのにスレッドが１６個しか出来ないの？と思われる方もいらっしゃるでしょうが、
実際にはそれより多いです。
上のプログラムでは１６より多い数のスレッドを起動していました。
スレッド16 ~ 31やそれ以降は、また別のまとめ（Coalescing）をしているということです。

[5回]

C#でOpenCL入門チュートリアル一覧

C#でOpenCLを使うシリーズの一覧をまとめます。
まだメモリフェンスだとか重要な機能を説明していないような気もしますが、どうもいじってみたところ、使わなくてもそれなりのことが出来るような気がします。
（メモリフェンスがなければ失敗するようなプログラムを書こうとしたのですが、仕様にない動作？により上手く動いてしまいました。環境によるのかもしれませんんが）

というわけでここで一区切り。
全部で１２の記事になりました：

感想

以前、似たような技術であるDirectX11のコンピュートシェーダーを扱いました。
それと比べるとOpenCLは簡単かつ柔軟な印象を受けました。

まずインターフェースが綺麗です。
DirectXではなんだか初期化が面倒くさいですがOpenCLはそんなことはありません。
まぁ「そんなことはありません」というと正確ではないかもしれませんが、グラフィックスは無視してGPGPUだけを考えているためか、ましなほうです。
また、スレッドグループのスレッド個数もDirectX11ではHLSLにハードコーディングしなくてはいけませんが、OpenCLではCPU側から柔軟に変更できます（意味があるかはあまりわかりませんが、柔軟といえば柔軟です）。

というふうにOpenCLのいいところだけ書きましたが、APIがC言語で書かれているので、C#との相互運用がちょっとめんどくさいですね。
誰かC#用OpenCLラッパーライブラリを書いてくれればだいぶ使いやすくなりそうです。

APIもラッパーを書きやすいタイプな気がします。
Managed Wrapper for OpenCLなんてのが出たらいいんですけどね。

[1回]

C#でOpenCL入門チュートリアルその１２　スレッドとグループの個数

GPGPUでは普通複数のスレッドを使います。
複数のスレッドで仕事をバーっとやってもらうのです。

そのスレッドの個数はCPU側で設定します。
ときに、GPU側でもスレッドの個数がいくつかを知りたいことがあります。

それには、get_global_size()関数を使います。
この関数は、今のコマンドで、スレッドが全部でいくつ実行されているかを取得します。
参考

size_t get_global_size(uint dimIndex)

dimIndexは、次元を表すインデックスです。X軸方向のスレッド個数を取得したいなら０，Y軸方向なら１，Zなら２です。

サンプルコード

MyProgram.cs

using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[6]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(2, 3), new MultiDimension(1, 1));

        var readBack = new float[6];
        commandQueue.ReadBuffer(buffer, readBack);

        foreach (var number in readBack)
        {
            Console.WriteLine(number);
        }
    }

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}

OpenCLWrappers.cs

using System;
using System.Runtime.InteropServices;
using System.Linq;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

    public void ReadBuffer<T>(Buffer buffer, T[] systemBuffer) where T : struct
    {
        GCHandle handle = GCHandle.Alloc(systemBuffer, GCHandleType.Pinned);

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        int error = OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );

        if (error != 0)
        {
            int paramValueSize = 0;
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                0, 
                null,
                out paramValueSize
                );
            System.Text.StringBuilder text = new System.Text.StringBuilder(paramValueSize);
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                paramValueSize,
                text,
                out paramValueSize);
            throw new Exception(text.ToString());
        }
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }
}

OpenCLFunctions.cs

using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );

    [DllImport("OpenCL.dll")]
    public static extern int clGetProgramBuildInfo(
        IntPtr program, 
        IntPtr device, 
        ProgramBuildInfoString paramName,
        int paramValueSize,
        System.Text.StringBuilder paramValue,
        out int paramValueSizeReturn
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }

    public MultiDimension(int x, int y)
    {
        X = x;
        Y = y;
        Z = 0;
        Dimension = 2;
    }
}

enum ProgramBuildInfoString
{
    Options = 0x1182,
    Log = 0x1183
}

myKernelProgram.cl

 
__kernel void myKernelFunction(__global float* items)
{
	items[get_global_id(0)] = 
		get_global_size(0);
		//get_global_size(1);
}

このプログラムは、スレッドを（２×３）個作ります。
そして、バッファの先頭に、０番目の次元のスレッド個数を書き込みます。
ここでは０番目の次元を書き込んでいるので、２という数字（２×３の２）を書き込みます。

実行結果はこうなります。

2
2
0
0
0
0

myKernelProgram.clのコメントアウト位置を変えると、今度は１番目の次元の個数を書き込みます。そうすると、バッファの先頭には３が書きこまれます（２×３の３です）。

3
3
0
0
0
0

グループのサイズ

以上では、スレッドの総数をGPU側で取得しました。
しかし、１グループ内のスレッド数を取得したい場合もあります。
それにはget_local_size関数を使います。
参考

size_t get_local_size(uint dimIndex)

dimIndexは取得するサイズの次元です。

たとえばスレッドを｛｛スレッド１，スレッド２、スレッド３｝、｛スレッド４、スレッド５,スレッド６｝｝と実行したとします。
全部で６つ。
１グループ３スレッドです。
そいう言う場合は、get_local_size(0)が３になります。
１グループ内のスレッド数を取得するのです。
サンプルを次のように改変します：

MyProgram.cs

using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[24]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(4, 6), new MultiDimension(2, 3));

        var readBack = new float[24];
        commandQueue.ReadBuffer(buffer, readBack);

        foreach (var number in readBack)
        {
            Console.WriteLine(number);
        }
    }

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}

myKernelFunction.cl

 
__kernel void myKernelFunction(__global float* items)
{
	items[get_global_id(0)] = 
		get_local_size(0);
		//get_local_size(1);
}

他のソースファイルは同じです。
これを実行すると、こうなります：

2
2
2
2
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0

このプログラムはまず４×６このスレッドを作り、２×３のスレッドを１グループにして、グループ分けしています。
そしてバッファの先頭に、ローカルサイズ（１グループのスレッド数）を書き込みます。
この場合は１グループにX方向には２ある（２×３の２）ので、２が書きこまれています。

グループの個数

スレッドが次のように実行されているとしましょう：
｛｛スレッド0、スレッド1｝、｛スレッド2、スレッド3｝、｛スレッド4、スレッド5｝｝
このとき、グループの個数は３個です。

グループの個数を得るには、get_num_groups()関数を使います。

参考

size_t get_num_groups(uint dimIndex)

dimIndexは次元のインデックスです。X軸方向のグループの数を得たい時には０。Y軸方向なら１。Zなら２です。

この関数の動作サンプルは次のようになります：

MyProgram.cs

using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[24]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(2, 12), new MultiDimension(2, 3));

        var readBack = new float[24];
        commandQueue.ReadBuffer(buffer, readBack);

        foreach (var number in readBack)
        {
            Console.WriteLine(number);
        }
    }

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}

myKernelProgram.cl

__kernel void myKernelFunction(__global float* items)
{
	items[get_global_id(0)] = 
		get_num_groups(0);
		//get_num_groups(1);
}

このプログラムは、２×１２個のスレッドを作り、それを２×３のスレッドを持つグループでグループ分けしています。
そうすると、グループの数は１×４となります。
結果、このプログラムは次のような文字列を出力します：

1

1
0

0
0

グループの数は１×４なので、その最初の１を出力したのです。

もしmyKernelProgram.clのコメントアウト位置を変えれば、４を出力するでしょう。

[0回]