-2

Cudaで画像全体の値の平均を計算したい。2D 配列のリダクションがどのように機能するかをテストするために、このカーネルを以下に記述します。最終出力 o は、すべての画像値の合計になります。入力 g は、すべてのピクセルに値 1 を持つ 2D 配列です。しかし、このプログラムの結果は合計として 0 です。私には少し奇妙です。

このチュートリアルhttp://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdfで 1D 配列の縮小を模倣します。この 2D フォームを書きます。私はCudaが初めてです。そして、潜在的なバグや改善への提案は大歓迎です!

コメントを 1 つ追加するだけです。1D配列で平均を計算するだけでも意味があることはわかっています。しかし、もっと活用して、より複雑なリダクション動作をテストしたいと思っています。それは正しくないかもしれません。しかし、ただのテストです。削減の一般的な慣行について、誰かが私にもっと提案してくれることを願っています。

#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

cudaEvent_t start, stop;
float elapsedTime;

__global__ void 
reduce(float *g, float *o, const int dimx, const int dimy)
{
extern __shared__ float sdata[];

unsigned int tid_x = threadIdx.x;
unsigned int tid_y = threadIdx.y;

unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int j = blockDim.y * blockIdx.y + threadIdx.y; 

if (i >= dimx || j >= dimy)
    return;

sdata[tid_x*blockDim.y + tid_y] = g[i*dimy + j];

__syncthreads();

for(unsigned int s_y = blockDim.y/2; s_y > 0; s_y >>= 1)
{
    if (tid_y < s_y)
    {
        sdata[tid_x * dimy + tid_y] += sdata[tid_x * dimy + tid_y + s_y];
    }
    __syncthreads();
}

for(unsigned int s_x = blockDim.x/2; s_x > 0; s_x >>= 1 )
{

    if(tid_x < s_x)
    {
        sdata[tid_x * dimy] += sdata[(tid_x + s_x) * dimy];
    }
    __syncthreads();
}

float sum;

if( tid_x == 0 && tid_y == 0)
{ 
    sum = sdata[0];
    atomicAdd (o, sum);   // The result should be the sum of all pixel values. But the program produce 0
}

//if(tid_x==0 && tid__y == 0 ) 
    //o[blockIdx.x] = sdata[0];
}

int
main()
{   
int dimx = 320;
int dimy = 160;
int num_bytes = dimx*dimy*sizeof(float);

float *d_a, *h_a, // device and host pointers
            *d_o=0, *h_o=0;

h_a = (float*)malloc(num_bytes);
h_o = (float*)malloc(sizeof(float));

srand(time(NULL));


for (int i=0; i < dimx; i++)
{   
    for (int j=0; j < dimy; j++)
    {
        h_a[i*dimy + j] = 1;
    }
}

cudaMalloc( (void**)&d_a, num_bytes );
cudaMalloc( (void**)&d_o, sizeof(int) );

cudaMemcpy( d_a, h_a, num_bytes, cudaMemcpyHostToDevice);
cudaMemcpy( d_o, h_o, sizeof(int), cudaMemcpyHostToDevice); 

dim3 grid, block;
block.x = 4;
block.y = 4;
grid.x = dimx / block.x;
grid.y = dimy / block.y;

cudaEventCreate(&start);
cudaEventRecord(start, 0);

int sizeofSharedMemory = dimx*dimy*sizeof(float);

reduce<<<grid, block, sizeofSharedMemory>>> (d_a, d_o, block.x, block.y);

cudaEventCreate(&stop);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

cudaEventElapsedTime(&elapsedTime, start, stop);
std::cout << "This kernel runs: " << elapsedTime << "ms" << std::endl; 

std::cout << block.x << " " << block.y << std::endl;
std::cout << grid.x << " " << grid.y << std::endl;
std::cout << dimx <<  " " << dimy << " " << dimx*dimy << std::endl;

cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );

std::cout << "The sum is:" << *h_o << std::endl;

free(h_a);
free(h_o);
cudaFree(d_a);
cudaFree(d_o);

}
4

1 に答える 1

11

基本的なcuda エラー チェックを行うと、reduce カーネルが実行されていないことがわかります。その理由は次のとおりです。

int dimx = 320;
int dimy = 160;
...
int sizeofSharedMemory = dimx*dimy*sizeof(float); // = 204800

reduce<<<grid, block, sizeofSharedMemory>>> (d_a, d_o, block.x, block.y);
                          ^
                          |
                         204800 is illegal here

204800 バイトの共有メモリを動的に (または他の方法で) 要求することはできません。最大値は 48K バイト弱です。

適切な cuda エラー チェックを行っていれば、カーネルが実行されていないことがわかり、起動構成 ( <<< ... >>> の間の数字) が無効であることを示唆する有益なエラー メッセージが表示されます。共有メモリはブロックごとに要求されます。各ブロックが 4x4 スレッド配列のみで構成されている場合、2D データ セット全体をカバーするのに十分な共有メモリを要求する必要があるのは賢明ではありません。おそらく、各 4x4 スレッド配列によってアクセスされるものに十分なデータが必要です。

cuda エラー チェックを使用してコードを適切に計測し、すべてのエラーを検出して修正したら、cuda-memcheck. これにより、追加レベルのエラー チェックが行われ、カーネル アクセス エラーが指摘されます。cuda-memcheck不特定の起動エラーが発生している場合にも使用でき、問題の特定に役立つ場合があります。

これらの基本的なトラブルシューティング手順を実行したら、他の人に助けを求めるのが理にかなっているかもしれません. ただし、最初に与えられたツールの力を使用してください。

また、戻ってきてこのコードを再度投稿して助けを求める前に、もう 1 つのエラーを指摘したいと思います。

これは役に立ちません:

std::cout << "The sum is:" << *h_o << std::endl;

cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );

デバイスからホストに合計をコピーする前に、合計を印刷しています。これらの手順の順序を逆にします。

cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );

std::cout << "The sum is:" << *h_o << std::endl;
于 2013-07-20T13:50:20.650 に答える