From 92cfa2be3a29cfdd8bf8ffc7fec76221ba761657 Mon Sep 17 00:00:00 2001 From: Yibing Liu Date: Sun, 17 Jun 2018 01:44:36 -0700 Subject: [PATCH] Avoid using dynamic array in cuda kernel --- paddle/fluid/operators/argsort_op.cc | 4 ++-- paddle/fluid/operators/argsort_op.cu | 7 +++---- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/paddle/fluid/operators/argsort_op.cc b/paddle/fluid/operators/argsort_op.cc index 8a44fd12ce0..ca9a884b983 100644 --- a/paddle/fluid/operators/argsort_op.cc +++ b/paddle/fluid/operators/argsort_op.cc @@ -38,8 +38,8 @@ class ArgsortOp : public framework::OperatorWithKernel { "dimension %d.", axis, num_dims); PADDLE_ENFORCE(in_dims.size() + axis >= 0, - "Attr(axis) %d of ArgsortOp plus the number of Input(X)'s " - "dimensions %d must be nonnegative.", + "Attr(axis) %d of ArgsortOp plus the rank %d of Input(X) " + "must be nonnegative.", axis, in_dims.size()); ctx->SetOutputDim("Out", in_dims); diff --git a/paddle/fluid/operators/argsort_op.cu b/paddle/fluid/operators/argsort_op.cu index 55ad4ce340d..fc64e51b34d 100644 --- a/paddle/fluid/operators/argsort_op.cu +++ b/paddle/fluid/operators/argsort_op.cu @@ -31,8 +31,9 @@ __global__ void ComputeTargetIdx(const int64_t* in_dims, int dims_size, int64_t* med_ids) { int64_t index = threadIdx.x + blockDim.x * blockIdx.x; if (index < n) { - int64_t* shape_out_axis = new int64_t[dims_size - 1]; - int64_t* dims_out_axis = new int64_t[dims_size - 1]; + const int max_rank = 9; // Max rank of a tensor allow in Fluid + int64_t shape_out_axis[max_rank - 1] = {0}; + int64_t dims_out_axis[max_rank - 1] = {0}; int64_t tmp = index; int64_t pos_in_axis = 0; int64_t i = dims_size - 2; @@ -57,8 +58,6 @@ __global__ void ComputeTargetIdx(const int64_t* in_dims, int dims_size, int64_t traget_idx = group * dim_axis + pos_in_axis; trg_idx[index] = traget_idx; med_ids[traget_idx] = pos_in_axis; - delete[] shape_out_axis; - delete[] dims_out_axis; } } -- GitLab