Skip to content

Commit

Permalink
fix format
Browse files Browse the repository at this point in the history
  • Loading branch information
zlsh80826 committed Apr 28, 2021
1 parent 757129b commit 4fd3118
Showing 1 changed file with 3 additions and 3 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -226,14 +226,13 @@ nvinfer1::DataType QkvToContextPluginDynamic::getOutputDataType(
}

template <typename T>
__global__ void apply_scale(T* data, T scale, int n) {
__global__ void apply_scale(T *data, T scale, int n) {
#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
int tid = blockIdx.x * blockDim.x + threadIdx.x;
data[tid] = data[tid] * scale;
#endif
}


int QkvToContextPluginDynamic::enqueue(
const nvinfer1::PluginTensorDesc *input_desc,
const nvinfer1::PluginTensorDesc *output_desc, const void *const *inputs,
Expand Down Expand Up @@ -304,7 +303,8 @@ int QkvToContextPluginDynamic::enqueue(
constexpr int threads = 128;
int blocks = (n_q + threads - 1) / threads;

apply_scale<<<blocks, threads, 0, stream>>>(tptr, static_cast<half>(scale_), n_q);
apply_scale<<<blocks, threads, 0, stream>>>(tptr, static_cast<half>(scale_),
n_q);

const platform::CUDADeviceContext &dev_ctx = *device_ctx;
operators::math::MultiHeadGPUComputeFunctor<half> multihead_compute_func;
Expand Down

1 comment on commit 4fd3118

@paddle-bot-old
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Congratulation! Your pull request passed all required CI. You could ask reviewer(s) to approve and merge. 🎉

Please sign in to comment.