Home
last modified time | relevance | path

Searched refs:HS_SLAB_THREADS (Results 1 – 19 of 19) sorted by relevance

/external/skqp/src/compute/hs/cl/intel/gen8/u64/
Dhs_kernels.cl569 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0) = r1;
570 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1) = r16;
571 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 2) = r2;
572 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 3) = r15;
573 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 4) = r3;
574 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 5) = r14;
575 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 6) = r4;
576 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 7) = r13;
577 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 8) = r5;
578 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 9) = r12;
[all …]
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/cl/intel/gen8/u32/
Dhs_kernels.cl375 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0) = r1;
376 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1) = r8;
377 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 2) = r2;
378 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 3) = r7;
379 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 4) = r3;
380 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 5) = r6;
381 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 6) = r4;
382 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 7) = r5;
415 r1 = HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0);
416 r8 = HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1);
[all …]
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/cuda/sm_35/u32/
Dhs_cuda_u32.cu1016 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0) = r1;
1017 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1) = r16;
1018 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 2) = r2;
1019 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 3) = r15;
1020 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 4) = r3;
1021 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 5) = r14;
1022 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 6) = r4;
1023 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 7) = r13;
1024 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 8) = r5;
1025 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 9) = r12;
[all …]
Dhs_cuda_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/cuda/sm_35/u64/
Dhs_cuda_u64.cu510 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0) = r1;
511 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1) = r8;
512 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 2) = r2;
513 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 3) = r7;
514 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 4) = r3;
515 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 5) = r6;
516 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 6) = r4;
517 HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 7) = r5;
550 r1 = HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 0);
551 r8 = HS_BX_LOCAL_V(2 * HS_SLAB_THREADS * 1);
[all …]
Dhs_cuda_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/
Dhs_glsl_macros.h48 HS_GLSL_WORKGROUP_SIZE(HS_SLAB_THREADS*slab_count,1,1); \
55 HS_GLSL_WORKGROUP_SIZE(HS_SLAB_THREADS*slab_count,1,1); \
61 HS_GLSL_WORKGROUP_SIZE(HS_SLAB_THREADS,1,1); \
67 HS_GLSL_WORKGROUP_SIZE(HS_SLAB_THREADS,1,1); \
73 HS_GLSL_WORKGROUP_SIZE(HS_SLAB_THREADS,1,1); \
116 (gl_GlobalInvocationID.x & ~(HS_SLAB_THREADS-1)) * \
118 (gl_LocalInvocationID.x & (HS_SLAB_THREADS-1))
121 extent[gmem_idx + HS_SLAB_THREADS * row_idx]
124 vout[gmem_idx + HS_SLAB_THREADS * row_idx] = reg
149 HS_SUBGROUP_ID() * (HS_SLAB_THREADS * slab_count) + \
[all …]
/external/skqp/src/compute/hs/cl/intel/
Dhs_cl_macros.h35 __attribute__((intel_reqd_sub_group_size(HS_SLAB_THREADS)))
44 __attribute__((reqd_work_group_size(HS_SLAB_THREADS*slab_count,1,1))) \
52 __attribute__((reqd_work_group_size(HS_SLAB_THREADS*slab_count,1,1))) \
103 (get_global_id(0) & ~(HS_SLAB_THREADS-1)) * HS_SLAB_HEIGHT + \
104 (get_local_id(0) & (HS_SLAB_THREADS-1))
107 extent[gmem_idx + HS_SLAB_THREADS * row_idx]
110 vout[gmem_idx + HS_SLAB_THREADS * row_idx] = reg
135 get_sub_group_id() * (HS_SLAB_THREADS * slab_count) + \
138 (get_sub_group_id() ^ 1) * (HS_SLAB_THREADS * slab_count) + \
139 (get_sub_group_local_id() ^ (HS_SLAB_THREADS - 1))
[all …]
/external/skqp/src/compute/hs/cuda/sm_35/
Dhs_cuda_macros.h91 HS_LAUNCH_BOUNDS(HS_SLAB_THREADS*slab_count,1) \
103 HS_LAUNCH_BOUNDS(HS_SLAB_THREADS*slab_count,HS_BS_SLABS/(1<<slab_count_ru_log2)) \
116 HS_LAUNCH_BOUNDS(HS_SLAB_THREADS*slab_count,HS_BS_SLABS/(1<<slab_count_log2)) \
158 HS_LAUNCH_BOUNDS(HS_SLAB_THREADS,1) \
193 (HS_GLOBAL_ID_X() & ~(HS_SLAB_THREADS-1)) * \
198 ((slab_offset + HS_GLOBAL_ID_X()) & ~(HS_SLAB_THREADS-1)) * \
202 extent[gmem_idx + HS_SLAB_THREADS * row_idx]
205 vout[gmem_idx + HS_SLAB_THREADS * row_idx] = reg
230 HS_WARP_ID_X() * (HS_SLAB_THREADS * slab_count) + \
233 (HS_WARP_ID_X() ^ 1) * (HS_SLAB_THREADS * slab_count) + \
[all …]
/external/skqp/src/compute/hs/vk/intel/
Dhs_glsl_macros.h32 const uint hs_subgroup_id = gl_LocalInvocationID.x / HS_SLAB_THREADS; \
33 const uint hs_subgroup_lane_id = gl_LocalInvocationID.x & (HS_SLAB_THREADS-1);
/external/skqp/src/compute/hs/cuda/
Dhs_cuda.inl584 <<<state->bx_ru,HS_SLAB_THREADS,0,state->streams[0]>>>
608 uint32_t const threads = HS_SLAB_THREADS << clean_slabs_log2;
666 <<<grid,HS_SLAB_THREADS * HS_HM_BLOCK_HEIGHT,0,stream>>>
764 <<<grid,HS_SLAB_THREADS * HS_FM_BLOCK_HEIGHT,0,stream>>>
777 <<<grid,HS_SLAB_THREADS * HS_FM_BLOCK_HEIGHT,0,stream>>>
810 <<<full_bs,HS_BS_SLABS*HS_SLAB_THREADS,0,stream>>>
821 <<<1,frac_bs*HS_SLAB_THREADS,0,stream>>>
822 (state->vout,state->vin,full_bs*HS_BS_SLABS*HS_SLAB_THREADS);
/external/skqp/src/compute/hs/vk/intel/gen8/u32/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/nvidia/sm_35/u64/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/amd/gcn/u64/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/intel/gen8/u64/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/nvidia/sm_35/u32/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro
/external/skqp/src/compute/hs/vk/amd/gcn/u32/
Dhs_config.h12 #define HS_SLAB_THREADS (1 << HS_SLAB_THREADS_LOG2) macro