From 8f1b4e0962857d402f9d017fcbf387ef0eceb7c4 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Thu, 01 Sep 2016 23:48:41 +0000
Subject: [PATCH] updates and things
---
src/blas_kernels.cu | 15 +++++++++++++++
1 files changed, 15 insertions(+), 0 deletions(-)
diff --git a/src/blas_kernels.cu b/src/blas_kernels.cu
index 3f7f1f9..271f017 100644
--- a/src/blas_kernels.cu
+++ b/src/blas_kernels.cu
@@ -368,6 +368,14 @@
if(i < N) X[i*INCX] = min(ALPHA, max(-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)
{
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -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