File: replace_nan_op.cu

package info (click to toggle)
pytorch 1.13.1%2Bdfsg-4
  • links: PTS, VCS
  • area: main
  • in suites: bookworm
  • size: 139,252 kB
  • sloc: cpp: 1,100,274; python: 706,454; ansic: 83,052; asm: 7,618; java: 3,273; sh: 2,841; javascript: 612; makefile: 323; xml: 269; ruby: 185; yacc: 144; objc: 68; lex: 44
file content (35 lines) | stat: -rw-r--r-- 779 bytes parent folder | download
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
#include "caffe2/core/context_gpu.h"
#include "caffe2/operators/replace_nan_op.h"

namespace caffe2 {

namespace {
template <typename T>
__global__ void
replace_nan_kernel(const T value, const int64_t size, const T* X, T* Y) {
  CUDA_1D_KERNEL_LOOP(i, size) {
    if (isnan(X[i])) {
      Y[i] = value;
    } else {
      Y[i] = X[i];
    }
  }
}
} // namespace

template <>
template <typename T>
void ReplaceNaNOp<CUDAContext>::ReplaceNaN(
    const T& value,
    const int64_t size,
    const T* X,
    T* Y) {
  replace_nan_kernel<<<
      CAFFE_GET_BLOCKS(size),
      CAFFE_CUDA_NUM_THREADS,
      0,
      context_.cuda_stream()>>>(value, size, X, Y);
  C10_CUDA_KERNEL_LAUNCH_CHECK();
}
REGISTER_CUDA_OPERATOR(ReplaceNaN, ReplaceNaNOp<CUDAContext>);
} // namespace caffe2