Home
last modified time | relevance | path

Searched refs:block_dim (Results 1 – 25 of 41) sorted by relevance

12

/external/tensorflow/tensorflow/core/kernels/
Dspacetobatch_functor_gpu.cu.cc67 for (int block_dim = NUM_BLOCK_DIMS - 1; block_dim >= 0; --block_dim) { in S2B() local
69 if (block_dim > 0) { in S2B()
70 offset %= args.block_shape[block_dim]; in S2B()
73 batch_tensor_pos[block_dim + 1] * args.block_shape[block_dim] + in S2B()
74 offset - args.pad_start[block_dim]; in S2B()
76 space_tensor_pos >= args.space_tensor_spatial_shape[block_dim]) { in S2B()
84 space_tensor_stride *= args.space_tensor_spatial_shape[block_dim]; in S2B()
85 if (block_dim == 0) { in S2B()
95 remaining_block_idx /= args.block_shape[block_dim]; in S2B()
117 for (int block_dim = 0; block_dim < NUM_BLOCK_DIMS; ++block_dim) { in operator ()() local
[all …]
Dspacetobatch_functor.cc107 for (int block_dim = 0; block_dim < NUM_BLOCK_DIMS; ++block_dim) { in operator ()() local
108 pad_start[block_dim] = paddings_tensor[block_dim * 2]; in operator ()()
109 block_shape[block_dim] = block_shape_tensor[block_dim]; in operator ()()
110 space_tensor_shape[block_dim] = space_tensor.dimension(block_dim + 1); in operator ()()
111 batch_tensor_shape[block_dim] = batch_tensor.dimension(block_dim + 1); in operator ()()
135 for (int block_dim = NUM_BLOCK_DIMS - 1; block_dim >= 0; --block_dim) { in operator ()() local
137 block_offsets[block_dim] = in operator ()()
138 block_dim > 0 ? block_index % block_shape[block_dim] : block_index; in operator ()()
139 block_index /= block_shape[block_dim]; in operator ()()
Dspacetobatch_op.cc101 for (int block_dim = 0; block_dim < block_dims; ++block_dim) { in SpaceToBatchOpCompute() local
102 block_shape_product *= block_shape[block_dim]; in SpaceToBatchOpCompute()
138 for (int block_dim = 0; block_dim < removed_prefix_block_dims; ++block_dim) { in SpaceToBatchOpCompute() local
139 const int64 size = orig_input_tensor.dim_size(block_dim + 1); in SpaceToBatchOpCompute()
146 for (int block_dim = removed_prefix_block_dims; in SpaceToBatchOpCompute() local
147 block_dim < block_dims - removed_suffix_block_dims; ++block_dim) { in SpaceToBatchOpCompute()
148 const int64 pad_start = paddings[2 * block_dim], in SpaceToBatchOpCompute()
149 pad_end = paddings[2 * block_dim + 1]; in SpaceToBatchOpCompute()
153 const int64 input_size = orig_input_tensor.dim_size(block_dim + 1); in SpaceToBatchOpCompute()
154 const int64 block_shape_value = block_shape[block_dim]; in SpaceToBatchOpCompute()
[all …]
Dbatchtospace_op.cc98 for (int block_dim = 0; block_dim < block_dims; ++block_dim) { in BatchToSpaceOpCompute() local
99 block_shape_product *= block_shape[block_dim]; in BatchToSpaceOpCompute()
140 for (int block_dim = 0; block_dim < removed_prefix_block_dims; ++block_dim) { in BatchToSpaceOpCompute() local
141 const int64 size = orig_input_tensor.dim_size(block_dim + 1); in BatchToSpaceOpCompute()
148 for (int block_dim = removed_prefix_block_dims; in BatchToSpaceOpCompute() local
149 block_dim < block_dims - removed_suffix_block_dims; ++block_dim) { in BatchToSpaceOpCompute()
150 const int64 crop_start = crops[2 * block_dim], in BatchToSpaceOpCompute()
151 crop_end = crops[2 * block_dim + 1]; in BatchToSpaceOpCompute()
154 const int64 input_size = orig_input_tensor.dim_size(block_dim + 1); in BatchToSpaceOpCompute()
155 const int64 block_shape_value = block_shape[block_dim]; in BatchToSpaceOpCompute()
[all …]
Dspacetobatch_benchmark_test.cc32 for (int block_dim = 0; block_dim < num_block_dims; ++block_dim) { in ConstructSpaceToBatchGraph() local
33 paddings_eigen_tensor(block_dim, 0) = paddings[block_dim].first; in ConstructSpaceToBatchGraph()
34 paddings_eigen_tensor(block_dim, 1) = paddings[block_dim].second; in ConstructSpaceToBatchGraph()
Ddepthwise_conv_op_gpu.h630 dim3 block_dim;
635 block_dim = dim3(kBlockDepth, args.in_cols, block_height);
644 block_dim = dim3(args.in_cols, block_height, kBlockDepth);
666 block_dim.x * block_dim.y * block_dim.z);
667 TF_CHECK_OK(GpuLaunchKernel(kernel, config.block_count, block_dim,
1598 dim3 block_dim;
1603 block_dim = dim3(kBlockDepth, args.in_cols, block_height);
1610 block_dim = dim3(args.in_cols, block_height, kBlockDepth);
1623 block_dim.x * block_dim.y * block_dim.z);
1624 TF_CHECK_OK(GpuLaunchKernel(kernel, config.block_count, block_dim,
/external/tensorflow/tensorflow/core/profiler/convert/
Dxplane_to_kernel_stats_db_test.cc91 EXPECT_EQ(kernel.block_dim().at(0), 1); in TEST()
92 EXPECT_EQ(kernel.block_dim().at(1), 1); in TEST()
93 EXPECT_EQ(kernel.block_dim().at(2), 1); in TEST()
109 EXPECT_EQ(kernel.block_dim().at(0), 32); in TEST()
110 EXPECT_EQ(kernel.block_dim().at(1), 1); in TEST()
111 EXPECT_EQ(kernel.block_dim().at(2), 1); in TEST()
127 EXPECT_EQ(kernel.block_dim().at(0), 64); in TEST()
128 EXPECT_EQ(kernel.block_dim().at(1), 1); in TEST()
129 EXPECT_EQ(kernel.block_dim().at(2), 1); in TEST()
/external/tensorflow/tensorflow/core/profiler/utils/
Dkernel_stats_utils_test.cc89 EXPECT_EQ(kernel.block_dim()[0], 32); in TEST()
90 EXPECT_EQ(kernel.block_dim()[1], 8); in TEST()
91 EXPECT_EQ(kernel.block_dim()[2], 4); in TEST()
110 EXPECT_EQ(kernel.block_dim()[0], 6); in TEST()
111 EXPECT_EQ(kernel.block_dim()[1], 5); in TEST()
112 EXPECT_EQ(kernel.block_dim()[2], 4); in TEST()
117 EXPECT_EQ(kernel.block_dim()[0], 1); in TEST()
118 EXPECT_EQ(kernel.block_dim()[1], 2); in TEST()
119 EXPECT_EQ(kernel.block_dim()[2], 3); in TEST()
Dkernel_stats_utils.cc174 lhs.block_dim(0), in operator ()()
175 lhs.block_dim(1), in operator ()()
176 lhs.block_dim(2), in operator ()()
189 rhs.block_dim(0), in operator ()()
190 rhs.block_dim(1), in operator ()()
191 rhs.block_dim(2), in operator ()()
211 lhs.block_dim(0) == rhs.block_dim(0) && in operator ()()
212 lhs.block_dim(1) == rhs.block_dim(1) && in operator ()()
213 lhs.block_dim(2) == rhs.block_dim(2) && in operator ()()
Dkernel_stats_utils.h75 wrap.key->block_dim(0), in AbslHashValue()
76 wrap.key->block_dim(1), in AbslHashValue()
77 wrap.key->block_dim(2), in AbslHashValue()
/external/tensorflow/tensorflow/lite/tools/optimize/sparsity/
Dformat_converter.cc62 int block_dim = 0; in FormatConverter() local
68 if (block_dim < block_map.size() && block_map[block_dim] == i) { in FormatConverter()
69 blocked_shape_[i] = shape[i] / block_size[block_dim]; in FormatConverter()
70 block_dim++; in FormatConverter()
242 int block_dim = 0; in FormatConverter() local
247 if (block_dim < block_map_.size() && block_map_[block_dim] == i) { in FormatConverter()
248 int orig_dim = traversal_order_[original_rank + block_dim]; in FormatConverter()
249 block_size_[block_dim] = sparsity.dim_metadata[orig_dim].dense_size; in FormatConverter()
251 block_dim++; in FormatConverter()
/external/tensorflow/tensorflow/stream_executor/
Ddevice_description.cc74 const BlockDim &block_dim = block_dim_limit(); in ToMap() local
76 absl::StrCat(block_dim.x, ",", block_dim.y, ",", block_dim.z); in ToMap()
/external/llvm-project/mlir/test/Dialect/GPU/
Dpromotion.mlir17 // CHECK-DAG: %[[bdx:.*]] = "gpu.block_dim"() {dimension = "x"}
18 // CHECK-DAG: %[[bdy:.*]] = "gpu.block_dim"() {dimension = "y"}
19 // CHECK-DAG: %[[bdz:.*]] = "gpu.block_dim"() {dimension = "z"}
71 // CHECK-DAG: %[[bdx:.*]] = "gpu.block_dim"() {dimension = "x"}
72 // CHECK-DAG: %[[bdy:.*]] = "gpu.block_dim"() {dimension = "y"}
73 // CHECK-DAG: %[[bdz:.*]] = "gpu.block_dim"() {dimension = "z"}
Dops.mlir35 %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
36 %bDimY = "gpu.block_dim"() {dimension = "y"} : () -> (index)
37 %bDimZ = "gpu.block_dim"() {dimension = "z"} : () -> (index)
Dall-reduce.mlir19 // CHECK: [[VAL_11:%.*]] = "gpu.block_dim"() {dimension = "x"} : () -> index
21 // CHECK: [[VAL_13:%.*]] = "gpu.block_dim"() {dimension = "y"} : () -> index
23 // CHECK: [[VAL_15:%.*]] = "gpu.block_dim"() {dimension = "z"} : () -> index
Doutlining.mlir51 // CHECK-NEXT: %[[BDIM:.*]] = "gpu.block_dim"() {dimension = "x"} : () -> index
52 // CHECK-NEXT: = "gpu.block_dim"() {dimension = "y"} : () -> index
53 // CHECK-NEXT: = "gpu.block_dim"() {dimension = "z"} : () -> index
/external/llvm-project/mlir/test/mlir-rocm-runner/
Dgpu-to-hsaco.mlir6 %block_dim = dim %arg1, %c0 : memref<?xf32>
8 threads(%tx, %ty, %tz) in (%block_x = %block_dim, %block_y = %c1, %block_z = %c1) {
Dvecadd.mlir6 %block_dim = dim %arg0, %c0 : memref<?xf32>
8 threads(%tx, %ty, %tz) in (%block_x = %block_dim, %block_y = %c1, %block_z = %c1) {
/external/tensorflow/tensorflow/tools/android/test/jni/object_tracking/
Dflow_cache.h270 int block_dim = kNumCacheLevels; in BlockDimForCacheLevel() local
273 block_dim *= kCacheBranchFactor; in BlockDimForCacheLevel()
275 return block_dim; in BlockDimForCacheLevel()
/external/tensorflow/tensorflow/core/util/
Dgpu_kernel_helper.h100 Status GpuLaunchKernel(void (*function)(Ts...), dim3 grid_dim, dim3 block_dim, in GpuLaunchKernel() argument
110 auto result = cudaLaunchKernel(func_ptr, grid_dim, block_dim, arg_ptrs.data(), in GpuLaunchKernel()
116 hipLaunchKernelGGL(function, grid_dim, block_dim, shared_memory_size_bytes, in GpuLaunchKernel()
/external/llvm-project/mlir/test/Conversion/GPUToSPIRV/
Dload-store.mlir55 %9 = "gpu.block_dim"() {dimension = "x"} : () -> index
56 %10 = "gpu.block_dim"() {dimension = "y"} : () -> index
57 %11 = "gpu.block_dim"() {dimension = "z"} : () -> index
Dbuiltins.mlir92 %0 = "gpu.block_dim"() {dimension = "x"} : () -> index
114 %0 = "gpu.block_dim"() {dimension = "y"} : () -> index
136 %0 = "gpu.block_dim"() {dimension = "z"} : () -> index
/external/tensorflow/tensorflow/core/kernels/image/
Dnon_max_suppression_op.cu.cc289 dim3 block_dim, thread_block; in NmsGpu() local
292 block_dim.x = num_blocks; in NmsGpu()
293 block_dim.y = num_blocks; in NmsGpu()
294 block_dim.z = 1; in NmsGpu()
299 TF_CHECK_OK(GpuLaunchKernel(NMSKernel<true>, block_dim, thread_block, 0, in NmsGpu()
303 TF_CHECK_OK(GpuLaunchKernel(NMSKernel<false>, block_dim, thread_block, 0, in NmsGpu()
/external/tensorflow/tensorflow/core/profiler/protobuf/
Dkernel_stats.proto16 repeated uint32 block_dim = 5; field
/external/tensorflow/tensorflow/compiler/tests/
Dspacetobatch_op_test.py54 for block_dim, block_shape_value in enumerate(block_shape):
55 reduced_size = padded.shape[block_dim + 1] // block_shape_value

12