| 120 | |
| 121 | template <typename T> |
| 122 | __global__ void BiasGradNHWC_SharedAtomics(int32 nthreads, |
| 123 | const T* output_backprop, |
| 124 | T* bias_backprop, int32 bias_size) { |
| 125 | typedef typename AccumulatorType<T>::type AccT; |
| 126 | GPU_DYNAMIC_SHARED_MEM_DECL(8, char, s_buf); |
| 127 | AccT* s_data = reinterpret_cast<AccT*>(s_buf); |
| 128 | for (int32 index = threadIdx.x; index < bias_size; index += blockDim.x) { |
| 129 | s_data[index] = AccT(0); |
| 130 | } |
| 131 | __syncthreads(); |
| 132 | |
| 133 | for (int32 index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; |
| 134 | index += blockDim.x * gridDim.x) { |
| 135 | int32 bias_offset = index % bias_size; |
| 136 | GpuAtomicAdd(s_data + bias_offset, AccT(ldg(output_backprop + index))); |
| 137 | } |
| 138 | __syncthreads(); |
| 139 | |
| 140 | for (int32 index = threadIdx.x; index < bias_size; index += blockDim.x) { |
| 141 | GpuAtomicAdd(bias_backprop + index, T(s_data[index])); |
| 142 | } |
| 143 | } |
| 144 | |
| 145 | template <typename T> |
| 146 | __global__ void BiasGradNCHW_SharedAtomics(const T* output_backprop, |
nothing calls this directly
no test coverage detected