Minor fix

This commit is contained in:
AlexeyAB 2020-03-23 20:03:56 +03:00
parent 864d1062f8
commit bb5eb481d1
1 changed files with 6 additions and 3 deletions

View File

@ -482,8 +482,9 @@ __global__ void scal_add_kernel(int N, float ALPHA, float BETA, float *X, int IN
__global__ void fill_kernel(int N, float ALPHA, float *X, int INCX) __global__ void fill_kernel(int N, float ALPHA, float *X, int INCX)
{ {
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; const int index = blockIdx.x*blockDim.x + threadIdx.x;
if(i < N) X[i*INCX] = ALPHA; if (index >= N) return;
X[index*INCX] = ALPHA;
} }
__global__ void mask_kernel_new_api(int n, float *x, float mask_num, float *mask, float val) __global__ void mask_kernel_new_api(int n, float *x, float mask_num, float *mask, float val)
@ -810,7 +811,9 @@ extern "C" void supp_ongpu(int N, float ALPHA, float * X, int INCX)
extern "C" void fill_ongpu(int N, float ALPHA, float * X, int INCX) extern "C" void fill_ongpu(int N, float ALPHA, float * X, int INCX)
{ {
fill_kernel<<<cuda_gridsize(N), BLOCK, 0, get_cuda_stream()>>>(N, ALPHA, X, INCX); //fill_kernel<<<cuda_gridsize(N), BLOCK, 0, get_cuda_stream()>>>(N, ALPHA, X, INCX);
//CHECK_CUDA(cudaPeekAtLastError());
fill_kernel << <get_number_of_blocks(N, BLOCK), BLOCK, 0, get_cuda_stream() >> >(N, ALPHA, X, INCX);
CHECK_CUDA(cudaPeekAtLastError()); CHECK_CUDA(cudaPeekAtLastError());
} }