CUDA,day-2,统计直方图

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <iostream>
#include <stdlib.h>
#include <conio.h>
using namespace std;


__global__ void func1(const unsigned char const * d_hist_data,
unsigned int * const d_bin_data)
{
const unsigned int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
const unsigned int tid = idx + idy*blockDim.x*gridDim.x;
const unsigned int value = d_hist_data[tid];


__uAtomicAdd(&(d_bin_data[value]), 1);
}


__global__ void func2(const unsigned char const * d_hist_data,
unsigned int * const d_bin_data)
{
const unsigned int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
const unsigned int tid = idx + idy*blockDim.x*gridDim.x;
const unsigned int value_32 = d_hist_data[tid];


__uAtomicAdd(&(d_bin_data[((value_32 & 0X000000FF))]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X0000FF00) >> 8)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X00FF0000) >> 16)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0XFF000000) >> 24)]), 1);
}


__shared__ unsigned int d_bin_data_shared[256];
__global__ void func3(const unsigned char const * d_hist_data,
unsigned int * const d_bin_data)
{
const unsigned int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
const unsigned int tid = idx + idy*blockDim.x*gridDim.x;
d_bin_data_shared[threadIdx.x] = 0;
const unsigned int value_32 = d_hist_data[tid];
__syncthreads();
__uAtomicAdd(&(d_bin_data[((value_32 & 0X000000FF))]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X0000FF00) >> 8)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X00FF0000) >> 16)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0XFF000000) >> 24)]), 1);
__syncthreads();
__uAtomicAdd(&(d_bin_data[threadIdx.x]), d_bin_data_shared[threadIdx.x]);
}


__global__ void func4(const unsigned char const * d_hist_data,
unsigned int * const d_bin_data,
unsigned int N)
{
const unsigned int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
const unsigned int tid = idx + idy*blockDim.x*gridDim.x;
d_bin_data_shared[threadIdx.x] = 0;
__syncthreads();
for (unsigned int i = 0, tid_offset = 0; 1 < N; i++, tid_offset += 256)
{
const unsigned int value_32 = d_hist_data[tid + tid_offset];
__uAtomicAdd(&(d_bin_data[((value_32 & 0X000000FF))]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X0000FF00) >> 8)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0X00FF0000) >> 16)]), 1);
__uAtomicAdd(&(d_bin_data[((value_32 & 0XFF000000) >> 24)]), 1);
}
__syncthreads();
__uAtomicAdd(&(d_bin_data[threadIdx.x]), d_bin_data_shared[threadIdx.x]);
}


int main()
{
return 0;
}