Home
last modified time | relevance | path

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

12

/external/eigen/unsupported/Eigen/CXX11/src/Tensor/
DTensorContractionCuda.h67 const Index lhs_store_idx_base = threadIdx.y * 72 + threadIdx.x * 9 + threadIdx.z; in EigenContractionKernelInternal()
68 const Index rhs_store_idx_base = threadIdx.y * 72 + threadIdx.z * 8 + threadIdx.x; in EigenContractionKernelInternal()
97 const Index load_idx_vert = threadIdx.x + 8 * threadIdx.y; in EigenContractionKernelInternal()
121 const Index lhs_horiz_0 = base_k + threadIdx.z + 0 * 8; \ in EigenContractionKernelInternal()
122 const Index lhs_horiz_1 = base_k + threadIdx.z + 1 * 8; \ in EigenContractionKernelInternal()
123 const Index lhs_horiz_2 = base_k + threadIdx.z + 2 * 8; \ in EigenContractionKernelInternal()
124 const Index lhs_horiz_3 = base_k + threadIdx.z + 3 * 8; \ in EigenContractionKernelInternal()
125 const Index lhs_horiz_4 = base_k + threadIdx.z + 4 * 8; \ in EigenContractionKernelInternal()
126 const Index lhs_horiz_5 = base_k + threadIdx.z + 5 * 8; \ in EigenContractionKernelInternal()
127 const Index lhs_horiz_6 = base_k + threadIdx.z + 6 * 8; \ in EigenContractionKernelInternal()
[all …]
DTensorReductionCuda.h115 const Index thread_id = blockIdx.x * blockDim.x + threadIdx.x; in ReductionInitKernel()
129 const Index first_index = blockIdx.x * BlockSize * NumPerThread + threadIdx.x; in FullReductionKernel()
136 if (threadIdx.x == 0) { in FullReductionKernel()
174 if ((threadIdx.x & (warpSize - 1)) == 0) { in FullReductionKernel()
178 if (gridDim.x > 1 && threadIdx.x == 0) { in FullReductionKernel()
205 const Index thread_id = blockIdx.x * blockDim.x + threadIdx.x; in ReductionInitKernelHalfFloat()
223 const Index first_index = blockIdx.x * BlockSize * NumPerThread + 2*threadIdx.x; in FullReductionKernelHalfFloat()
250 if ((threadIdx.x & (warpSize - 1)) == 0) { in FullReductionKernelHalfFloat()
265 eigen_assert(threadIdx.x == 1); in ReductionCleanupKernelHalfFloat()
390 const Index thread_id = blockIdx.x * blockDim.x + threadIdx.x;
[all …]
DTensorConvolution.h587 for (int p = first_plane + threadIdx.y; p < numPlanes; p += plane_stride) {
590 const int plane_kernel_offset = threadIdx.y * num_x_input;
592 for (int i = threadIdx.x; i < num_x_input; i += blockDim.x) {
603 for (int i = threadIdx.x; i < num_x_output; i += blockDim.x) {
641 for (int p = first_plane + threadIdx.z; p < numPlanes; p += plane_stride) {
644 const int plane_kernel_offset = threadIdx.z * num_y_input;
648 for (int j = threadIdx.y; j < num_y_input; j += blockDim.y) {
651 for (int i = threadIdx.x; i < num_x_input; i += blockDim.x) {
663 for (int j = threadIdx.y; j < num_y_output; j += blockDim.y) {
665 for (int i = threadIdx.x; i < num_x_output; i += blockDim.x) {
[all …]
DTensorRandom.h22 assert(threadIdx.z == 0); in get_random_seed()
24 blockIdx.x * blockDim.x + threadIdx.x + in get_random_seed()
25 gridDim.x * blockDim.x * (blockIdx.y * blockDim.y + threadIdx.y); in get_random_seed()
/external/clang/test/SemaCUDA/
Dcuda-builtin-vars.cu7 out[i++] = threadIdx.x; in kernel()
8 threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}} in kernel()
9 out[i++] = threadIdx.y; in kernel()
10 threadIdx.y = 0; // expected-error {{no setter defined for property 'y'}} in kernel()
11 out[i++] = threadIdx.z; in kernel()
12 threadIdx.z = 0; // expected-error {{no setter defined for property 'z'}} in kernel()
43 …__cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of cla… in kernel()
46 …threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_thre… in kernel()
49 …void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_thr… in kernel()
/external/clang/test/CodeGenCUDA/
Dcuda-builtin-vars.cu9 out[i++] = threadIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x() in kernel()
10 out[i++] = threadIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y() in kernel()
11 out[i++] = threadIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z() in kernel()
/external/tensorflow/tensorflow/core/kernels/
Dbias_op_gpu.cu.cc130 for (int32 index = threadIdx.x; index < bias_size; index += blockDim.x) { in BiasGradNHWC_SharedAtomics()
135 for (int32 index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; in BiasGradNHWC_SharedAtomics()
142 for (int32 index = threadIdx.x; index < bias_size; index += blockDim.x) { in BiasGradNHWC_SharedAtomics()
156 for (int32 index = threadIdx.x; index < kSDataSize; index += blockDim.x) { in BiasGradNCHW_SharedAtomics()
167 for (int32 index = group_index * blockDim.x + threadIdx.x; in BiasGradNCHW_SharedAtomics()
178 int bias_offset = threadIdx.x % 32; in BiasGradNCHW_SharedAtomics()
184 int32 thread_index = threadIdx.x; in BiasGradNCHW_SharedAtomics()
Dconcat_lib_gpu_impl.cu.cc41 IntType gidx = blockIdx.x * blockDim.x + threadIdx.x; in concat_fixed_kernel()
44 IntType gidy = blockIdx.y * blockDim.y + threadIdx.y; in concat_fixed_kernel()
69 IntType gidx = blockIdx.x * blockDim.x + threadIdx.x; in concat_variable_kernel()
77 IntType lidx = threadIdx.y * blockDim.x + threadIdx.x; in concat_variable_kernel()
108 IntType gidy = blockIdx.y * blockDim.y + threadIdx.y; in concat_variable_kernel()
Dreduction_gpu_kernels.cu.h198 const int tid = threadIdx.x;
238 int warp_index = threadIdx.x / 32;
240 const int lane = threadIdx.x % 32;
243 int gid = threadIdx.x + blockIdx.x * blockDim.x;
304 const int lane = threadIdx.x % 32;
308 rows_per_warp * (blockIdx.y * blockDim.y + threadIdx.y);
336 sum, static_cast<int>(threadIdx.x + i * num_cols), 0xffffffff);
340 if (lane < num_cols) partial_sums[lane * 33 + threadIdx.y] = sum;
344 if (threadIdx.y == 0 && threadIdx.x < num_cols) {
345 value_type s = partial_sums[threadIdx.x * 33];
[all …]
Dsplit_lib_gpu.cu.cc124 IntType gidx = blockIdx.x * blockDim.x + threadIdx.x; in split_v_kernel()
132 IntType lidx = threadIdx.y * blockDim.x + threadIdx.x; in split_v_kernel()
163 IntType gidy = blockIdx.y * blockDim.y + threadIdx.y; in split_v_kernel()
Dbucketize_op_gpu.cu.cc46 int32 lidx = threadIdx.y * blockDim.x + threadIdx.x; in BucketizeCustomKernel()
Drandom_op_gpu.h135 const int32 thread_id = blockIdx.x * blockDim.x + threadIdx.x;
174 const int32 thread_id = blockIdx.x * blockDim.x + threadIdx.x;
Ddepthwise_conv_op_gpu.h218 const int thread_depth = threadIdx.x; in DepthwiseConv2dGPUKernelNHWCSmall()
219 const int thread_col = threadIdx.y; in DepthwiseConv2dGPUKernelNHWCSmall()
220 const int thread_row = threadIdx.z; in DepthwiseConv2dGPUKernelNHWCSmall()
501 const int thread_col = threadIdx.x; in DepthwiseConv2dGPUKernelNCHWSmall()
502 const int thread_row = threadIdx.y; in DepthwiseConv2dGPUKernelNCHWSmall()
503 const int thread_depth = threadIdx.z; in DepthwiseConv2dGPUKernelNCHWSmall()
1214 const int thread_depth = threadIdx.x;
1215 const int thread_col = threadIdx.y;
1216 const int thread_row = threadIdx.z;
1485 const int thread_col = threadIdx.x;
[all …]
Dcheck_numerics_op_gpu.cu.cc40 const int32 thread_id = blockIdx.x * blockDim.x + threadIdx.x; in CheckNumericsKernel()
/external/skia/src/compute/hs/cuda/sm_35/
Dhs_cuda_macros.h182 #define HS_GLOBAL_ID_X() (blockDim.x * blockIdx.x + threadIdx.x)
183 #define HS_LOCAL_ID_X() threadIdx.x
184 #define HS_WARP_ID_X() (threadIdx.x / 32)
185 #define HS_LANE_ID() (threadIdx.x & 31)
/external/skqp/src/compute/hs/cuda/sm_35/
Dhs_cuda_macros.h182 #define HS_GLOBAL_ID_X() (blockDim.x * blockIdx.x + threadIdx.x)
183 #define HS_LOCAL_ID_X() threadIdx.x
184 #define HS_WARP_ID_X() (threadIdx.x / 32)
185 #define HS_LANE_ID() (threadIdx.x & 31)
/external/llvm/test/Analysis/DivergenceAnalysis/NVPTX/
Ddiverge.ll6 ; return (n < 0 ? a + threadIdx.x : b + threadIdx.x)
26 ; if (threadIdx.x < 5) // divergent: data dependent
45 ; if (threadIdx.x >= 5) { // divergent
48 ; // c here is divergent because it is sync dependent on threadIdx.x >= 5
/external/swiftshader/third_party/llvm-7.0/llvm/test/Analysis/DivergenceAnalysis/NVPTX/
Ddiverge.ll6 ; return (n < 0 ? a + threadIdx.x : b + threadIdx.x)
26 ; if (threadIdx.x < 5) // divergent: data dependent
45 ; if (threadIdx.x >= 5) { // divergent
48 ; // c here is divergent because it is sync dependent on threadIdx.x >= 5
/external/eigen/test/
Dcuda_common.h11 dim3 threadIdx, blockDim, blockIdx; variable
26 int i = threadIdx.x + blockIdx.x*blockDim.x; in run_on_cuda_meta_kernel()
/external/tensorflow/tensorflow/contrib/rnn/kernels/
Dlstm_ops_gpu.cu.cc91 const int batch_id = blockIdx.x * blockDim.x + threadIdx.x; in lstm_gates()
92 const int act_id = blockIdx.y * blockDim.y + threadIdx.y; in lstm_gates()
206 const int gid = blockDim.x * blockIdx.x + threadIdx.x; in concat_xh()
304 const int batch_id = blockIdx.x * blockDim.x + threadIdx.x; in lstm_gates_bprop()
305 const int act_id = blockIdx.y * blockDim.y + threadIdx.y; in lstm_gates_bprop()
/external/tensorflow/tensorflow/examples/adding_an_op/
Dcuda_op_kernel.cu.cc22 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; in AddOneKernel()
/external/tensorflow/tensorflow/tools/ci_build/builds/user_ops/
Dcuda_op_kernel.cu.cc22 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; in AddOneKernel()
/external/tensorflow/tensorflow/compiler/xla/service/gpu/llvm_gpu_backend/tests_data/
Dsaxpy.ll9 @threadIdx = external addrspace(1) global %struct.uint3
38 …truct.uint3, %struct.uint3* addrspacecast (%struct.uint3 addrspace(1)* @threadIdx to %struct.uint3…
90 …truct.uint3, %struct.uint3* addrspacecast (%struct.uint3 addrspace(1)* @threadIdx to %struct.uint3…
/external/tensorflow/tensorflow/contrib/tensorrt/custom_plugin_examples/
Dinc_op_kernel.cu.cc33 int i = blockDim.x * blockIdx.x + threadIdx.x; in VecInc()
/external/clang/lib/Headers/
Dcuda_builtin_vars.h112 __CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx; variable

12