1 
__kernel void kmeans_kernel(__global float* data, int points, 
                            __global float* centroids, int clusters, 
                             int dimensions)
{           

    //extern __shared__ float storage_space[];
     __local float storage_space[];

    __local int iterations;
    __local float *means; 
    __local float *index; 
    __local float *mindist;
    __local float *s_data;


    iterations = points / ( get_global_size(0)) + 1;    


    if( get_local_id(0) == 0 ){
        s_data[get_local_id(0)] = data[get_local_id(0)];
        means=&storage_space[0];
        index=&storage_space[get_local_size(0)];
        mindist=&storage_space[2*get_local_size(0)];
    }

    //data = &data[blockDim.x * blockIdx.x + threadIdx.x];
    data = &data[get_global_id(0)];         

    while( iterations )
    {
        mindist[get_local_id(0)] = 3.402823466e+38F;                
        index[get_local_id(0)] = 0;             

        for( short j = 0; j < clusters; j++ )
        {                       
            if( get_local_id(0) <= dimensions )
                //means[get_local_id(0)] = centroids[get_local_id(0)+j*c_pitch];
                means[get_local_id(0)] = centroids[get_local_id(0)+j];
            //__syncthreads();      
            barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

            if( !(data[get_local_id(0)] - s_data[get_local_id(0)] > points - 1) )
            {
                float dist = distance_gpu_transpose( means, data, dimensions);
                if( dist < mindist[get_local_id(0)] )
                {
                    mindist[get_local_id(0)] = dist;
                    index[get_local_id(0)] = j;
                }                   
            }           
            //__syncthreads();          
            barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
        }

        if( !(data[get_local_id(0)] - s_data[get_local_id(0)] > points - 1) )                               
                data[0] = index[get_local_id(0)];

        data += (get_global_size(0));               
        if( get_local_id(0) == 0 )
            --iterations;
        //__syncthreads();
        barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
    }
}

ATI 3200 グラフィックス カードで AMD プロセッサを使用しています。これは openCL をサポートしていませんが、残りのコードは CPU 自体で正常に動作しています。

今回のコードの問題は、私にとって非常に複雑です。カーネルが実行された後、 を使用してデバイス メモリ変数を読み取ることができませんclEnqueueReadBuffer。デバッグ中に、この時点で壊れており、次のように述べています。

Unhandled exception at 0x10001098 in CL_kmeans.exe: 0xC000001D: Illegal Instruction.

ここでブレークを押すと、

No symbols are loaded for any call stack frame. The source code cannot be displayed.

ここで何が問題になる可能性がありますか？解決策を教えてください。私のカーネルコードは上記のとおりであり、データを読み取るために使用するステートメントは、

ret = clEnqueueReadBuffer(command_queue, gpu_data, CL_TRUE, 0,sizeof( float ) * instances->cols* 1 , instances->data, 0, NULL, NULL);

ここで考えられるエラーを確認するにはどうすればよいですか?

4

0 に答える 0