From 0f645836f193e75c4c3b718369e6fab15b5d19c5 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Wed, 11 Feb 2015 03:41:03 +0000
Subject: [PATCH] Detection is back, baby\!
---
src/im2col_kernels.cu | 14 +++++++-------
1 files changed, 7 insertions(+), 7 deletions(-)
diff --git a/src/im2col_kernels.cu b/src/im2col_kernels.cu
index feaf44d..a82c2dc 100644
--- a/src/im2col_kernels.cu
+++ b/src/im2col_kernels.cu
@@ -3,7 +3,7 @@
#include "cuda.h"
}
-__global__ void im2col_pad_kernel(float *im, int offset,
+__global__ void im2col_pad_kernel(float *im,
int channels, int height, int width,
int ksize, int stride, float *data_col)
{
@@ -32,13 +32,13 @@
int im_row = h_offset + h * stride - pad;
int im_col = w_offset + w * stride - pad;
- int im_index = offset + im_col + width*(im_row + height*im_channel);
+ int im_index = im_col + width*(im_row + height*im_channel);
float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];
data_col[col_index] = val;
}
-__global__ void im2col_nopad_kernel(float *im, int offset,
+__global__ void im2col_nopad_kernel(float *im,
int channels, int height, int width,
int ksize, int stride, float *data_col)
{
@@ -65,13 +65,13 @@
int im_row = h_offset + h * stride;
int im_col = w_offset + w * stride;
- int im_index = offset + im_col + width*(im_row + height*im_channel);
+ int im_index = im_col + width*(im_row + height*im_channel);
float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];
data_col[col_index] = val;
}
-extern "C" void im2col_ongpu(float *im, int offset,
+extern "C" void im2col_ongpu(float *im,
int channels, int height, int width,
int ksize, int stride, int pad, float *data_col)
{
@@ -87,7 +87,7 @@
size_t n = channels_col*height_col*width_col;
- if(pad)im2col_pad_kernel<<<cuda_gridsize(n),BLOCK>>>(im, offset, channels, height, width, ksize, stride, data_col);
- else im2col_nopad_kernel<<<cuda_gridsize(n),BLOCK>>>(im, offset, channels, height, width, ksize, stride, data_col);
+ if(pad)im2col_pad_kernel<<<cuda_gridsize(n),BLOCK>>>(im, channels, height, width, ksize, stride, data_col);
+ else im2col_nopad_kernel<<<cuda_gridsize(n),BLOCK>>>(im, channels, height, width, ksize, stride, data_col);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0