| | |
| | | extern "C" { |
| | | #include "blas.h" |
| | | #include "cuda.h" |
| | | #include "utils.h" |
| | | } |
| | | |
| | | __global__ void axpy_kernel(int N, float ALPHA, float *X, int OFFX, int INCX, float *Y, int OFFY, int INCY) |
| | |
| | | if(i < N) X[i*INCX] *= ALPHA; |
| | | } |
| | | |
| | | __global__ void mask_kernel(int n, float *x, float *mask) |
| | | __global__ void mask_kernel(int n, float *x, float mask_num, float *mask) |
| | | { |
| | | int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; |
| | | if(i < n && mask[i] == 0) x[i] = 0; |
| | | if(i < n && mask[i] == mask_num) x[i] = mask_num; |
| | | } |
| | | |
| | | __global__ void copy_kernel(int N, float *X, int OFFX, int INCX, float *Y, int OFFY, int INCY) |
| | |
| | | check_error(cudaPeekAtLastError()); |
| | | } |
| | | |
| | | extern "C" void mask_ongpu(int N, float * X, float * mask) |
| | | extern "C" void mask_ongpu(int N, float * X, float mask_num, float * mask) |
| | | { |
| | | mask_kernel<<<cuda_gridsize(N), BLOCK>>>(N, X, mask); |
| | | mask_kernel<<<cuda_gridsize(N), BLOCK>>>(N, X, mask_num, mask); |
| | | check_error(cudaPeekAtLastError()); |
| | | } |
| | | |