Home
last modified time | relevance | path

Searched refs:input_size (Results 1 – 25 of 386) sorted by relevance

12345678910>>...16

/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/
Dcast_impl.cu97 __global__ void CastKernel(const int input_size, const S *input_addr, T *output_addr) { in CastKernel() argument
98 …for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < input_size; pos += blockDim.x * gri… in CastKernel()
104 void Cast(const int input_size, const S *input_addr, T *output_addr, cudaStream_t stream) { in Cast() argument
105 …CastKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, input_addr, output_addr… in Cast()
108 template void Cast(const int input_size, const int8_t *input_addr, int8_t *output_addr, cudaStream_…
109 template void Cast(const int input_size, const int8_t *input_addr, int16_t *output_addr, cudaStream…
110 template void Cast(const int input_size, const int8_t *input_addr, int32_t *output_addr, cudaStream…
111 template void Cast(const int input_size, const int8_t *input_addr, int64_t *output_addr, cudaStream…
112 template void Cast(const int input_size, const int8_t *input_addr, uint8_t *output_addr, cudaStream…
113 template void Cast(const int input_size, const int8_t *input_addr, uint16_t *output_addr, cudaStrea…
[all …]
Dsmooth_l1_loss_impl.cu21 __global__ void SmoothL1LossKernel(const int input_size, const float beta, const T *prediction, con… in SmoothL1LossKernel() argument
23 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < input_size; i += blockDim.x * gridDim.x) { in SmoothL1LossKernel()
34 void SmoothL1Loss(const int &input_size, const float &beta, const T *prediction, const T *target, T… in SmoothL1Loss() argument
36 …SmoothL1LossKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, beta, predictio… in SmoothL1Loss()
40 __global__ void SmoothL1LossGradKernel(const int input_size, const float beta, const T *prediction,… in SmoothL1LossGradKernel() argument
42 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < input_size; i += blockDim.x * gridDim.x) { in SmoothL1LossGradKernel()
55 void SmoothL1LossGrad(const int &input_size, const float &beta, const T *prediction, const T *targe… in SmoothL1LossGrad() argument
57 …SmoothL1LossGradKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, beta, predi… in SmoothL1LossGrad()
61 template void SmoothL1Loss<float>(const int &input_size, const float &beta, const float *prediction,
63 template void SmoothL1LossGrad<float>(const int &input_size, const float &beta, const float *predic…
Dloss_with_reduction_impl.cu30 __global__ void Copy(T *loss, T *tmp_loss, int reduction, int input_size) { in Copy() argument
33 loss[0] /= castT(loss[0], input_size); in Copy()
137 __global__ void KLDivLossKernel(const int input_size, const int reduction, const T *input_x, const … in KLDivLossKernel() argument
141 … for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < input_size; i += blockDim.x * gridDim.x) { in KLDivLossKernel()
147 … for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < input_size; i += blockDim.x * gridDim.x) { in KLDivLossKernel()
156 void KLDivLoss(const int &input_size, const int &reduction, const T *input_x, const T *input_y, T *… in KLDivLoss() argument
159 …KLDivLossKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, reduction, input_x… in KLDivLoss()
162 if (input_size % 2 == 1) { in KLDivLoss()
163 AddTile<<<1, 1, 0, stream>>>(tmp_loss, input_size - 1); in KLDivLoss()
165 for (int stride = input_size / 2; stride > 0; stride >>= 1) { in KLDivLoss()
[all …]
Dscatter_nd.cu22 …ScatterNdKernel(S *indices, T *update, T *output, const size_t block_size, const size_t input_size, in ScatterNdKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in ScatterNdKernel()
50 void ScatterNd(S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in ScatterNd() argument
53 …GET_BLOCKS(output_size), GET_THREADS, 0, stream>>>(indices, update, output, block_size, input_size, in ScatterNd()
60 … const size_t &input_size, const size_t &output_size, const size_t &indices_dim_0,
64 const size_t &input_size, const size_t &output_size,
68 … const size_t &input_size, const size_t &output_size, const size_t &indices_dim_0,
72 const size_t &input_size, const size_t &output_size,
76 … const size_t &input_size, const size_t &output_size, const size_t &indices_dim_0,
80 … const size_t &input_size, const size_t &output_size, const size_t &indices_dim_0,
[all …]
Dtensor_scatter_min.cu23 … const size_t input_size, const size_t output_size, const size_t indices_dim_0, in TensorScatterMinKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in TensorScatterMinKernel()
50 …Min(T *input, S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in TensorScatterMin() argument
54 …input, indices, update, output, block_size, input_size, output_size, indices_dim_0, indices_dim_1,… in TensorScatterMin()
61 const size_t &block_size, const size_t &input_size,
67 const size_t &block_size, const size_t &input_size,
73 const size_t &block_size, const size_t &input_size,
80 … const size_t &input_size, const size_t &output_size,
85 const size_t &block_size, const size_t &input_size,
92 const size_t &block_size, const size_t &input_size,
[all …]
Dtensor_scatter_max.cu23 … const size_t input_size, const size_t output_size, const size_t indices_dim_0, in TensorScatterMaxKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in TensorScatterMaxKernel()
50 …Max(T *input, S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in TensorScatterMax() argument
54 …input, indices, update, output, block_size, input_size, output_size, indices_dim_0, indices_dim_1,… in TensorScatterMax()
61 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
66 const size_t &block_size, const size_t &input_size,
72 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
78 … const size_t &input_size, const size_t &output_size,
83 const size_t &input_size, const size_t &output_size,
89 const size_t &block_size, const size_t &input_size,
[all …]
Dtensor_scatter_sub.cu23 … const size_t input_size, const size_t output_size, const size_t indices_dim_0, in TensorScatterSubKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in TensorScatterSubKernel()
50 …Sub(T *input, S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in TensorScatterSub() argument
54 …input, indices, update, output, block_size, input_size, output_size, indices_dim_0, indices_dim_1,… in TensorScatterSub()
61 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
66 const size_t &block_size, const size_t &input_size,
72 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
78 … const size_t &input_size, const size_t &output_size,
83 const size_t &input_size, const size_t &output_size,
89 const size_t &block_size, const size_t &input_size,
[all …]
Dtensor_scatter_update.cu23 … const size_t input_size, const size_t output_size, const size_t indices_dim_0, in TensorScatterUpdateKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in TensorScatterUpdateKernel()
50 …ate(T *input, S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in TensorScatterUpdate() argument
54 …input, indices, update, output, block_size, input_size, output_size, indices_dim_0, indices_dim_1,… in TensorScatterUpdate()
60 const size_t &block_size, const size_t &input_size,
65 const size_t &block_size, const size_t &input_size,
70 const size_t &block_size, const size_t &input_size,
75 const size_t &block_size, const size_t &input_size,
81 … const size_t &input_size, const size_t &output_size,
85 const size_t &block_size, const size_t &input_size,
[all …]
Dreverse_v2_impl.cu21 const int64_t* axis, size_t input_size, size_t axis_size) { in ReverseV2() argument
22 …for (int64_t gt_id = blockIdx.x * blockDim.x + threadIdx.x; gt_id < input_size; gt_id += blockDim.… in ReverseV2()
37 size_t input_size, size_t axis_size, cudaStream_t cuda_stream) { in CalReverseV2() argument
38 …ReverseV2<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input, output, input_shape, str… in CalReverseV2()
39 input_size, axis_size); in CalReverseV2()
44 … const int64_t* axis, size_t input_size, size_t axis_size, cudaStream_t cuda_stream);
47 … const int64_t* axis, size_t input_size, size_t axis_size, cudaStream_t cuda_stream);
50 … const int64_t* strides, const int64_t* axis, size_t input_size, size_t axis_size,
54 … const int64_t* strides, const int64_t* axis, size_t input_size, size_t axis_size,
58 … const int64_t* strides, const int64_t* axis, size_t input_size, size_t axis_size,
[all …]
Dtile_impl.cu20 __global__ void Tile(const size_t output_size, const size_t input_size, const size_t shape_size, in Tile() argument
39 pos_size = input_size; in Tile()
50 void CalTile(const size_t output_size, const size_t input_size, const size_t shape_size, const size… in CalTile() argument
52 …Tile<<<GET_BLOCKS(output_size), GET_THREADS, 0, cuda_stream>>>(output_size, input_size, shape_size… in CalTile()
57 template void CalTile<double>(const size_t output_size, const size_t input_size, const size_t shape…
60 template void CalTile<float>(const size_t output_size, const size_t input_size, const size_t shape_…
63 template void CalTile<half>(const size_t output_size, const size_t input_size, const size_t shape_s…
66 template void CalTile<int16_t>(const size_t output_size, const size_t input_size, const size_t shap…
69 template void CalTile<int>(const size_t output_size, const size_t input_size, const size_t shape_si…
72 template void CalTile<int64_t>(const size_t output_size, const size_t input_size, const size_t shap…
[all …]
Dbce_with_logits_loss_impl.cu92 void CalBCEWithLogitsLoss(const size_t input_size, const T *predict, const T *target, const size_t … in CalBCEWithLogitsLoss() argument
98 FillAndBroadcast<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>( in CalBCEWithLogitsLoss()
99 input_size, shape_size, pos_weight_shape, input_shape, pos_weight, shape_broadcasted); in CalBCEWithLogitsLoss()
101 …FillWithoutBroadcast<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input_size, pos_weig… in CalBCEWithLogitsLoss()
104 …BCEWithLogitsLossMain<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input_size, predict… in CalBCEWithLogitsLoss()
107 …FillAndBroadcast<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input_size, shape_size, … in CalBCEWithLogitsLoss()
110 … FillWithoutBroadcast<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input_size, weight, in CalBCEWithLogitsLoss()
113 …Mul<<<GET_BLOCKS(input_size), GET_THREADS, 0, cuda_stream>>>(input_size, shape_broadcasted, output… in CalBCEWithLogitsLoss()
117 template void CalBCEWithLogitsLoss<half>(const size_t input_size, const half *predict, const half *…
123 template void CalBCEWithLogitsLoss<float>(const size_t input_size, const float *predict, const floa…
Dl2_loss.cu22 __global__ void L2LossKernel(const size_t input_size, const T *input , T *output) { in L2LossKernel() argument
24 …for (size_t id = blockIdx.x * blockDim.x + threadIdx.x; id < input_size; id += blockDim.x * gridDi… in L2LossKernel()
37 void L2Loss(const size_t input_size, const T *input , T *output, cudaStream_t stream) { in L2Loss() argument
39 L2LossKernel<<<GET_BLOCKS(input_size), GET_THREADS, 0, stream>>>(input_size, input, output); in L2Loss()
42 template void L2Loss<float>(const size_t input_size, const float *input , float *output, cudaStream…
43 template void L2Loss<half>(const size_t input_size, const half *input , half *output, cudaStream_t …
Dtensor_scatter_add.cu23 … const size_t input_size, const size_t output_size, const size_t indices_dim_0, in TensorScatterAddKernel() argument
26 for (size_t read_index = blockIdx.x * blockDim.x + threadIdx.x; read_index < input_size; in TensorScatterAddKernel()
50 …Add(T *input, S *indices, T *update, T *output, const size_t &block_size, const size_t &input_size, in TensorScatterAdd() argument
54 …input, indices, update, output, block_size, input_size, output_size, indices_dim_0, indices_dim_1,… in TensorScatterAdd()
60 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
64 const size_t &block_size, const size_t &input_size,
69 const size_t &block_size, const size_t &input_size,
74 … const size_t &block_size, const size_t &input_size, const size_t &output_size,
79 … const size_t &input_size, const size_t &output_size,
83 const size_t &input_size, const size_t &output_size,
[all …]
/third_party/mindspore/mindspore/ccsrc/frontend/parallel/tensor_layout/
Dtensor_redistribution.cc242 Status TensorRedistribution::ComputePermuteCost(double input_size, const Shape &attrs) { in ComputePermuteCost() argument
249 forward_comm_cost_ += input_size * ALLTOALL_SCALE_FACTOR; in ComputePermuteCost()
250 backward_comm_cost_ += input_size * ALLTOALL_SCALE_FACTOR; in ComputePermuteCost()
251 comm_cost_ += COST_FACTOR * input_size * ALLTOALL_SCALE_FACTOR; in ComputePermuteCost()
255 computation_cost_ += input_size; in ComputePermuteCost()
256 memory_cost_ += input_size; in ComputePermuteCost()
260 computation_cost_ += (input_size + input_size * dev_num + input_size * dev_num); in ComputePermuteCost()
261 memory_cost_ += (input_size * dev_num + input_size * dev_num + input_size); in ComputePermuteCost()
266 Status TensorRedistribution::ComputeConcatCost(double input_size, const Shape &attrs) { in ComputeConcatCost() argument
275 forward_comm_cost_ += input_size * dev_num * ALLGATHER_REDUCESCATTER_SCALE_FACTOR; in ComputeConcatCost()
[all …]
/third_party/mindspore/mindspore/lite/test/ut/src/runtime/kernel/arm/fp32_grad/
Dactivation_grad_fp32_tests.cc42 size_t input_size; in TEST_F() local
44 …input_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(input_path.c_str(), &input_size)); in TEST_F()
46 EXPECT_EQ(input_size, output_data_size * sizeof(float)); in TEST_F()
49 auto yt_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(yt_path.c_str(), &input_size)); in TEST_F()
51 EXPECT_EQ(input_size, output_data_size * sizeof(float)); in TEST_F()
96 size_t input_size; in TEST_F() local
98 …input_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(input_path.c_str(), &input_size)); in TEST_F()
102 auto yt_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(yt_path.c_str(), &input_size)); in TEST_F()
147 size_t input_size; in TEST_F() local
149 …input_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(input_path.c_str(), &input_size)); in TEST_F()
[all …]
/third_party/mindspore/tests/st/ops/gpu/
Dtest_batchtospace_op.py29 input_size = 1
31 input_size = input_size*i
32 data_np = np.arange(input_size).reshape(input_shape).astype(nptype)
44 input_size = 1
46 input_size = input_size*i
59 input_size = 1
61 input_size = input_size*i
68 arr_input = Tensor(np.arange(input_size).reshape(input_shape).astype(nptype))
Dtest_spacetobatch_op.py29 input_size = 1
31 input_size = input_size*i
32 data_np = np.arange(input_size).reshape(input_shape).astype(nptype)
44 input_size = 1
46 input_size = input_size*i
63 input_size = 1
65 input_size = input_size*i
76 arr_input = Tensor(np.arange(input_size).reshape(input_shape).astype(nptype))
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/cpu/nnacl/fp32_grad/
Dbinary_cross_entropy.c20 static void BinaryCrossEntropyLossKernel(const int input_size, const int reduction, const float *in… in BinaryCrossEntropyLossKernel() argument
24 for (int i = 0; i < input_size; i++) { in BinaryCrossEntropyLossKernel()
30 for (int i = 0; i < input_size; i++) { in BinaryCrossEntropyLossKernel()
38 void BinaryCrossEntropy(const int input_size, const int reduction, const float *input_x, const floa… in BinaryCrossEntropy() argument
41 BinaryCrossEntropyLossKernel(input_size, reduction, input_x, input_y, weight, loss, tmp_loss); in BinaryCrossEntropy()
43 if (input_size % 2 == 1) { in BinaryCrossEntropy()
44 tmp_loss[0] += tmp_loss[input_size - 1]; in BinaryCrossEntropy()
46 for (int stride = input_size / 2; stride > 0; stride = stride / 2) { in BinaryCrossEntropy()
57 loss[0] /= input_size; in BinaryCrossEntropy()
/third_party/mindspore/mindspore/lite/test/ut/src/runtime/kernel/arm/fp16_grad/
Dactivation_grad_fp16_test.cc35 size_t input_size; in TEST_F() local
37 …input_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(input_path.c_str(), &input_size)); in TEST_F()
39 EXPECT_EQ(input_size, output_data_size * sizeof(float)); in TEST_F()
42 auto yt_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(yt_path.c_str(), &input_size)); in TEST_F()
44 EXPECT_EQ(input_size, output_data_size * sizeof(float)); in TEST_F()
47 …ata = reinterpret_cast<const float *>(mindspore::lite::ReadFile(output_path.c_str(), &input_size)); in TEST_F()
49 EXPECT_EQ(input_size, output_data_size * sizeof(float)); in TEST_F()
92 size_t input_size; in TEST_F() local
94 …input_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(input_path.c_str(), &input_size)); in TEST_F()
98 auto yt_data = reinterpret_cast<float *>(mindspore::lite::ReadFile(yt_path.c_str(), &input_size)); in TEST_F()
[all …]
/third_party/boost/libs/compute/include/boost/compute/algorithm/detail/
Dinplace_reduce.hpp37 size_t input_size = iterator_range_size(first, last); in inplace_reduce() local
38 if(input_size < 2){ in inplace_reduce()
46 size_t block_count = input_size / (block_size * values_per_thread); in inplace_reduce()
47 if(block_count * block_size * values_per_thread != input_size) in inplace_reduce()
100 while(input_size > 1){ in inplace_reduce()
102 kernel.set_arg(input_size_arg, static_cast<uint_>(input_size)); in inplace_reduce()
111 input_size = in inplace_reduce()
113 std::ceil(float(input_size) / (block_size * values_per_thread) in inplace_reduce()
117 block_count = input_size / (block_size * values_per_thread); in inplace_reduce()
118 if(block_count * block_size * values_per_thread != input_size) in inplace_reduce()
/third_party/boost/boost/compute/algorithm/detail/
Dinplace_reduce.hpp37 size_t input_size = iterator_range_size(first, last); in inplace_reduce() local
38 if(input_size < 2){ in inplace_reduce()
46 size_t block_count = input_size / (block_size * values_per_thread); in inplace_reduce()
47 if(block_count * block_size * values_per_thread != input_size) in inplace_reduce()
100 while(input_size > 1){ in inplace_reduce()
102 kernel.set_arg(input_size_arg, static_cast<uint_>(input_size)); in inplace_reduce()
111 input_size = in inplace_reduce()
113 std::ceil(float(input_size) / (block_size * values_per_thread) in inplace_reduce()
117 block_count = input_size / (block_size * values_per_thread); in inplace_reduce()
118 if(block_count * block_size * values_per_thread != input_size) in inplace_reduce()
/third_party/mindspore/tests/st/ops/ascend/
Dtest_gru_op.py40 …def __init__(self, input_size, hidden_size, num_layers, has_bias, batch_first, bidirectional, drop… argument
42 …self.gru = nn.GRU(input_size=input_size, hidden_size=hidden_size, num_layers=num_layers, has_bias=…
50 … def __init__(self, num_layers, has_bias, input_size, num_directions, hidden_size, bidirectional): argument
53 self.input_size = input_size
68 … layer_input_size = self.input_size if layer == 0 else self.hidden_size * self.num_directions
96 input_size = 32
103 … fact = GRUWeightBias(num_layers, has_bias, input_size, num_directions, hidden_size, bidirectional)
111 …net = GRU(input_size=input_size, hidden_size=16, num_layers=num_layers, has_bias=has_bias, batch_f…
121 …net_pynative = GRU(input_size=input_size, hidden_size=16, num_layers=num_layers, has_bias=has_bias,
137 input_size = 32
[all …]
Dtest_rnn_op.py40 …def __init__(self, input_size, hidden_size, num_layers, has_bias, batch_first, bidirectional, drop… argument
42 …self.rnn = nn.RNN(input_size=input_size, hidden_size=hidden_size, num_layers=num_layers, has_bias=…
50 … def __init__(self, num_layers, has_bias, input_size, num_directions, hidden_size, bidirectional): argument
53 self.input_size = input_size
68 … layer_input_size = self.input_size if layer == 0 else self.hidden_size * self.num_directions
96 input_size = 32
103 … fact = RNNWeightBias(num_layers, has_bias, input_size, num_directions, hidden_size, bidirectional)
111 …net = RNN(input_size=input_size, hidden_size=16, num_layers=num_layers, has_bias=has_bias, batch_f…
121 …net_pynative = RNN(input_size=input_size, hidden_size=16, num_layers=num_layers, has_bias=has_bias,
137 input_size = 32
[all …]
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/nn/
Drmsprop_gpu_kernel.h95 size_t input_size = size_ * sizeof(T); in InitSizeLists() local
97 input_size_list_.push_back(input_size); in InitSizeLists()
98 input_size_list_.push_back(input_size); in InitSizeLists()
99 input_size_list_.push_back(input_size); in InitSizeLists()
101 input_size_list_.push_back(input_size); in InitSizeLists()
102 output_size_list_.push_back(input_size); in InitSizeLists()
104 input_size_list_.push_back(input_size); in InitSizeLists()
105 input_size_list_.push_back(input_size); in InitSizeLists()
106 input_size_list_.push_back(input_size); in InitSizeLists()
107 input_size_list_.push_back(input_size); in InitSizeLists()
[all …]
/third_party/mbedtls/programs/psa/
Dcrypto_examples.c63 size_t input_size, in cipher_operation() argument
73 while( bytes_written != input_size ) in cipher_operation()
75 bytes_to_write = ( input_size - bytes_written > part_size ? in cipher_operation()
77 input_size - bytes_written ); in cipher_operation()
102 size_t input_size, in cipher_encrypt() argument
119 status = cipher_operation( &operation, input, input_size, part_size, in cipher_encrypt()
133 size_t input_size, in cipher_decrypt() argument
149 status = cipher_operation( &operation, input, input_size, part_size, in cipher_decrypt()
212 input_size = 100, in cipher_example_encrypt_decrypt_aes_cbc_pkcs7_multi() enumerator
222 uint8_t iv[block_size], input[input_size], in cipher_example_encrypt_decrypt_aes_cbc_pkcs7_multi()
[all …]

12345678910>>...16