OrgLion-ML
/
darknet
mirror of https://github.com/AlexeyAB/darknet.git
1
0
Fork 0
Code Issues Projects Releases Wiki Activity

Fix [implicit] layer

This commit is contained in:
AlexeyAB 2021-05-12 04:16:19 +03:00
parent 846c79b6d4
commit 81b768bae0
2 changed files with 4 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
src/blas_kernels.cu
View File

@ -2460,12 +2460,14 @@ __global__ void backward_implicit_kernel(int size, int batch, int nweights, floa
const int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; const int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if (id >= size) return; if (id >= size) return;
weight_updates_gpu[id % nweights] += delta_gpu[id]; for (int i = 0; i < batch; ++i) {
weight_updates_gpu[id] += delta_gpu[id + i * nweights];
}
} }
extern "C" void backward_implicit_gpu(int batch, int nweights, float *weight_updates_gpu, float *delta_gpu) extern "C" void backward_implicit_gpu(int batch, int nweights, float *weight_updates_gpu, float *delta_gpu)
{ {
int size = batch * nweights; int size = nweights;
backward_implicit_kernel << <cuda_gridsize(size), BLOCK, 0, get_cuda_stream() >> > (size, batch, nweights, weight_updates_gpu, delta_gpu); backward_implicit_kernel << <cuda_gridsize(size), BLOCK, 0, get_cuda_stream() >> > (size, batch, nweights, weight_updates_gpu, delta_gpu);
CHECK_CUDA(cudaPeekAtLastError()); CHECK_CUDA(cudaPeekAtLastError());
} }

1
src/representation_layer.c
View File

@ -67,7 +67,6 @@ void forward_implicit_layer(const layer l, network_state state)
void backward_implicit_layer(const layer l, network_state state) void backward_implicit_layer(const layer l, network_state state)
{ {
int i; int i;
#pragma omp parallel for
for (i = 0; i < l.nweights * l.batch; ++i) { for (i = 0; i < l.nweights * l.batch; ++i) {
l.weight_updates[i % l.nweights] += l.delta[i]; l.weight_updates[i % l.nweights] += l.delta[i];
} }