From ae53edc6a42833116b469bdbf288ae9adf01671a Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 06 Oct 2016 06:35:06 +0000
Subject: [PATCH] fix bug thing
---
src/blas_kernels.cu | 17 ++++++++++++++++-
1 files changed, 16 insertions(+), 1 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 3f7f1f9..0391e2e 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -365,7 +365,15 @@
__global__ void constrain_kernel(int N, float ALPHA, float *X, int INCX)
{
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
- if(i < N) X[i*INCX] = min(ALPHA, max(-ALPHA, X[i*INCX]));
+ if(i < N) X[i*INCX] = fminf(ALPHA, fmaxf(-ALPHA, X[i*INCX]));
+}
+
+__global__ void supp_kernel(int N, float ALPHA, float *X, int INCX)
+{
+ int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+ if(i < N) {
+ if((X[i*INCX] * X[i*INCX]) < (ALPHA * ALPHA)) X[i*INCX] = 0;
+ }
}
__global__ void scal_kernel(int N, float ALPHA, float *X, int INCX)
@@ -552,6 +560,12 @@
check_error(cudaPeekAtLastError());
}
+extern "C" void supp_ongpu(int N, float ALPHA, float * X, int INCX)
+{
+ supp_kernel<<<cuda_gridsize(N), BLOCK>>>(N, ALPHA, X, INCX);
+ check_error(cudaPeekAtLastError());
+}
+
extern "C" void fill_ongpu(int N, float ALPHA, float * X, int INCX)
{
fill_kernel<<<cuda_gridsize(N), BLOCK>>>(N, ALPHA, X, INCX);
@@ -633,6 +647,7 @@
}
+
__global__ void weighted_sum_kernel(int n, float *a, float *b, float *s, float *c)
{
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
--
Gitblit v1.10.0