Форум программистов [Powered by Invision Power Board]

Код

const UINT AVG2_THREADS_IN_BLOCK=256;
// кол-во задач в блоке должно быть степенью двойки.

__global__
void average_kernel_WORDS2(const WORD* const d_in_data, UINT data_size, float* const d_out_value) {
    // d_in_data - входные данные в формате 16-ти разрядных беззнаковых целых
    // data_size - размер входных данных в элементах
    __shared__ float avg[AVG2_THREADS_IN_BLOCK];
    
    unsigned int index = blockIdx.x*blockDim.x + threadIdx.x;
    // копируем данные в shared и делим их на data_size
    if (index<data_size) {
        avg[threadIdx.x]=float(d_in_data[index])/data_size;
    }
    else {
        avg[threadIdx.x]=0.;
    }
    __syncthreads();

        // считаем сумму в shared через scan. Результат в последней ячейке
    for (unsigned int i=2; i<=blockDim.x; i<<=1) {
        if ((threadIdx.x +1) % i == 0) {
            avg[threadIdx.x]+=avg[threadIdx.x-i/2];
        }
    }
    __syncthreads();

    if (threadIdx.x==blockDim.x-1) { // последняя задача выдаёт результат
        atomicAdd(d_out_value,float(avg[threadIdx.x]));
    }
}


// функция обёртка над kernel для вычисления среднего арифметического

bool CUDA_Average2(WORD* d_in_data, UINT data_size, float& result) {
        
    // определяем кол-во блоков.
    UINT blocksNum=data_size/AVG2_THREADS_IN_BLOCK;
    if (data_size%AVG2_THREADS_IN_BLOCK!=0) blocksNum++;
    if (blocksNum==0) blocksNum++;

    dim3 threads(AVG2_THREADS_IN_BLOCK);
    dim3 blocks(blocksNum);

    // выделяем память под массив промежуточных результатов
    float* d_out_value=0; float h_out_value=0;
    try {
        if (cudaMalloc(&d_out_value, sizeof(float)) !=  cudaSuccess) throw (1);
        if (cudaMemcpy(d_out_value, &h_out_value, sizeof(float), cudaMemcpyHostToDevice) !=  cudaSuccess) throw (1);
    
        average_kernel_WORDS2<<<blocks,threads,sizeof(float)*AVG2_THREADS_IN_BLOCK>>>(d_in_data,data_size,d_out_value);
        if (cudaMemcpy(&h_out_value, d_out_value, sizeof(float), cudaMemcpyDeviceToHost) !=  cudaSuccess) throw (1);
    }
    catch (...) {
        if (d_out_value!=0) cudaFree(d_out_value);
        return false;
    }

    result=h_out_value;
    cudaFree(d_out_value);

    return true;
}