Reply to comment

Submitted by blackswan on Thu, 11/19/2009 - 17:39.

я открыл код reduction из toolkit, однако такой код мне не очень понятен, я написал свой, видимо это тоже каскадное суммирование, он работает, однако мне кажется, что выдает неверные значения, так как есть реализация обычная для тех же данных (того же изображения), и эти ср. арифметические отличаются от тех. Я разбил задачу на две стадии, сделав предварительный массив сумм.

__global__ void __meanStage1(float *A, float *Sum, long N, long M)
{
    long iS = blockIdx.y * gridDim.x + blockIdx.x;
    for (long j= blockIdx.y * BLOCK_SIZE; j< (blockIdx.y+1) * BLOCK_SIZE ; j++)
                for (long i= blockIdx.x * BLOCK_SIZE; i< (blockIdx.x+1) * BLOCK_SIZE ; i++)
            if( i < N && j < M)
                Sum[iS] = Sum[iS] + A[j * M + i];
    __syncthreads();
}
__global__ void __meanStage2(float *a,float *mean, float *sum, long dX, long dY, long N, long M)
{
        *mean = 0;
    for (long i=0;i<dX * dY;i++)
        *mean = *mean + sum[i];
    if (M*N!=0)
    *mean = (*mean)/(N*M);
    else
    *mean = 0;
    __syncthreads();
}
void mean(float *a,float &Means, long N, long M)
{
        float *sum,*mean;
        long iS,jS;
        iS = N / BLOCK_SIZE;
        jS = M / BLOCK_SIZE;
        cudaMalloc((void**) &mean , sizeof(float));
        cudaMalloc((void**) &sum , iS * jS * sizeof(float));
        cudaMemset(sum,0, iS * jS * sizeof(float));
    dim3 dimBlock1(1, 1);
    dim3 dimGrid1(N/BLOCK_SIZE, M/BLOCK_SIZE);
    __meanStage1<<<dimGrid1, dimBlock1>>>(a, sum, N, M);
    dim3 dimBlock2(1, 1);
    dim3 dimGrid2(1, 1);
    __meanStage2<<<dimGrid2, dimBlock2>>>(a, mean, sum, iS, jS, N, M);
    cudaFree(sum);
    cudaMemcpy(&Means, mean, sizeof(float),cudaMemcpyDeviceToHost);
    cudaFree(mean);
}

GPGPU.ru

Использование видеокарт для вычислений

Reply to comment

Reply

Вход

Navigation

Cвежие комментарии

Новое на форуме

Популярно

Популярно сегодня