0

2 つのカーネルがあります (どちらも 1 回だけ実行されるため、この例では globalWorkSize は 1 です)。

最初のカーネル ( kernel_Calc()) は、いくつかの値を計算し、__globalメモリに格納します。この例では、変換行列を計算 (3D 空間でポイントを変換する変換行列を設定) し、origo を変換します。

inline 
float4 mul( const float4 M[ 4 ], const float4 v)
{
   float4 r;
   r.x = dot( v, M[ 0 ] );
   r.y = dot( v, M[ 1 ] );
   r.z = dot( v, M[ 2 ] );
   r.w = dot( v, M[ 3 ] );
   return r;
}

__kernel
void kernel_Calc( __global float4* g_TransformationMatrices, __global float3* g_Point3D )
{
    __private float4 transformationMatrix[ 4 ];

    transformationMatrix [ 0 ] = (float4) ( 1.0f, 0.0f, 0.0f, 0.0f );
    transformationMatrix [ 1 ] = (float4) ( 0.0f, 1.0f, 0.0f, 10.0f );
    transformationMatrix [ 2 ] = (float4) ( 0.0f, 0.0f, 1.0f, 0.0f );
    transformationMatrix [ 3 ] = (float4) ( 0.0f, 0.0f, 0.0f, 1.0f );

    g_TransformationMatrices[ 0 ] = transformationMatrix[ 0 ];
    g_TransformationMatrices[ 1 ] = transformationMatrix[ 1 ];
    g_TransformationMatrices[ 2 ] = transformationMatrix[ 2 ];
    g_TransformationMatrices[ 3 ] = transformationMatrix[ 3 ];


    float4 point4D = (float4) ( 0.0f, 0.0f, 0.0f, 1.0f );
    float4 point4DTransformed = mul( transformationMatrix, point4D);

    g_Point3D[ 0 ] = (float3) ( point4DTransformed.x / point4DTransformed.w ,
                                point4DTransformed.y / point4DTransformed.w ,
                                point4DTransformed.z / point4DTransformed.w );
}

ホスト側では、計算__globalされたバッファーを関数を使用して__constantバッファー (CL_MEM_READ_ONLYバッファー) にコピーしclEnqueueCopyBuffer()ます。__constant（メモリからの読み取りがメモリからの読み取りよりも高速であることを願っているため、これを行い__globalます。この機能を使用すると、デバイス側でバッファのコピーを実行できます。__globalホストにコピーしてからコピーする必要はありません__constant。）

2 番目のカーネル ( kernel_Test()) は、計算された値をホスト側で読み取ることができる__global変数 ( ) にロードしようとします。__global float4* testこれsizeStructは、整数のみを含むユーザー定義の構造体です (これは行列と変換されたポイントの数です)。__constant2 番目と 3 番目のパラメーターは、関数で満たされたメモリ内のバッファーclEnqueueCopyBuffer()です。

struct sizeStruct
{
    int m_Size;
};

__kernel 
void kernel_Test( __constant struct sizeStruct* c_SS,
                  __constant float4* c_TransformationMatrices,
                  __constant float3* c_Points3D,
                  __global float4 *test )
{                   
    test[ 0 ] = c_TransformationMatrices[ 0 ];
    test[ 1 ] = c_TransformationMatrices[ 1 ];
    test[ 2 ] = c_TransformationMatrices[ 2 ];
    test[ 3 ] = c_TransformationMatrices[ 3 ];
}

問題は、プログラムを実行すると、テスト変数に次のものが含まれることです。

1.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000

ただし、次のものが含まれている必要があります。

1.000000, 0.000000, 0.000000, 0.000000
0.000000, 1.000000, 0.000000, 10.000000
0.000000, 0.000000, 1.000000, 0.000000
0.000000, 0.000000, 0.000000, 1.000000

__constant正しいデータが含まれている変数を (ホスト メモリにコピーして)チェックしました。コードは、私のプログラムの簡易版です。これが、不要な操作やパラメーターが含まれている可能性がある理由です。この例はテスト済みで、説明したとおりに機能します。

__constant float3* c_Points3Dカーネル パラメーターをカーネル パラメーターに変更すると__global float3* c_Points3D(ただし、関数でいっぱいになった read_only バッファーを引き続き使用しclEnqueueCopyBuffer()ます)、正常に動作します。パラメータを削除しても機能し__constant struct sizeStruct* c_SSます。したがって、kernel_Test の引数のアドレス空間に何か問題があるようですが、問題は__constant-> __globalcopy で発生します。

nvidia geforce gtx 690 でプログラムを実行していますが、デバイス (およびプラットフォーム) を intel i7-3930k (intel sdk を使用) に変更できます。intel-i7 CPU を使用すると、カーネル コードを変更しなくてもすべて正常に動作します。

Q1:この奇妙な動作が表示されるのはなぜですか? 誰かが私が間違っていることを知っていますか?

Q2:アドレス空間修飾子cl_mem_read_onlyを使用してバッファーを作成し、使用することは合法ですか?__global

4

1 に答える 1