| | |
| | * Copyright (c) Facebook, Inc. and its affiliates. |
| | * |
| | * This source code is licensed under the MIT license found in the |
| | * LICENSE file in the root directory of this source tree. |
| | */ |
| |
|
| | #include <ATen/ATen.h> |
| | #include <c10/cuda/CUDAStream.h> |
| |
|
| | #include <cuda.h> |
| | #include <cuda_fp16.h> |
| | #include <cuda_runtime.h> |
| |
|
| | #include <algorithm> |
| | #include <functional> |
| | #include <iostream> |
| | #include <stdexcept> |
| | #include <utility> |
| | #include <vector> |
| |
|
| | #include <stdlib.h> |
| | #include <assert.h> |
| | #include <math.h> |
| |
|
| | #define SHFL_MASK 0xffffffff |
| |
|
| | template<int FS, int SB, int padding_l, typename scalar_t> |
| | __global__ |
| | void dynamicconv_forward_kernel(const scalar_t* input, |
| | const scalar_t* weight, |
| | int minibatch, |
| | int sequenceLength, |
| | int numFeatures, |
| | int numFiltersInBlock, |
| | int numHeads, |
| | scalar_t* output); |
| |
|
| | template<int FS, int SB, int padding_l, typename scalar_t> |
| | __global__ |
| | void dynamicconv_backward_kernel( |
| | const scalar_t* gradOutput, |
| | const scalar_t* input, |
| | const scalar_t* weight, |
| | int minibatch, |
| | int sequenceLength, |
| | int numFeatures, |
| | int numFiltersInBlock, |
| | int numHeads, |
| | scalar_t* gradWeight, |
| | scalar_t* gradInput); |
| |
|
