 }
 
 template <typename T>
-__global__ void RenormKernelFunc4(T* x_data, T* out_data, int64_t size,
+__global__ void RenormKernelFunc4(const T* x_data, T* out_data, int64_t size,
                                   T* dim_value, int64_t dimension_each,
                                   int64_t dim_divisor) {
   int64_t i = ((int64_t)blockIdx.x) * blockDim.x + threadIdx.x;

// See the License for the specific language governing permissions and
// limitations under the License.

#include <algorithm>
#include "paddle/fluid/framework/op_registry.h"
#include "paddle/fluid/operators/elementwise/elementwise_op_impl.cu.h"
#include "paddle/fluid/operators/reduce_ops/reduce_functor_op.h"
#include "paddle/fluid/operators/reduce_ops/reduce_op.cu.h"
#include "paddle/fluid/operators/reduce_ops/reduce_op.h"
#include "paddle/fluid/operators/renorm_op.h"
#include "paddle/fluid/operators/utils.h"
#include "paddle/fluid/platform/device/gpu/gpu_primitives.h"
#include "stdio.h"

namespace paddle {
namespace operators {

                        ops::CUDARenormKernel<double>);

REGISTER_OP_CUDA_KERNEL(renorm_grad, ops::CUDAGradRenormKernel<float>,
                        ops::CUDAGradRenormKernel<double>);
