/aosp_15_r20/external/mesa3d/src/amd/vulkan/radix_sort/shaders/ |
H A D | prefix.h | 142 if (gl_LocalInvocationID.x < RS_SWEEP_0_SIZE) // subgroup has inactive invocations in rs_prefix() 145 const uint32_t h0_red = RS_PREFIX_SWEEP0(gl_LocalInvocationID.x); in rs_prefix() 148 RS_PREFIX_SWEEP0(gl_LocalInvocationID.x) = h0_inc - h0_red; in rs_prefix() 159 if (gl_LocalInvocationID.x < RS_SWEEP_0_SIZE) // 32 invocations in rs_prefix() 161 const uint32_t h0_red = RS_PREFIX_SWEEP0(gl_LocalInvocationID.x); in rs_prefix() 164 RS_PREFIX_SWEEP0(gl_LocalInvocationID.x) = h0_inc - h0_red; in rs_prefix() 176 const uint32_t idx0 = (ii * RS_WORKGROUP_SIZE) + gl_LocalInvocationID.x; in rs_prefix() 193 if (gl_LocalInvocationID.x < RS_SWEEP_1_SIZE) // 4 invocations in rs_prefix() 195 const uint32_t h1_red = RS_PREFIX_SWEEP1(gl_LocalInvocationID.x); in rs_prefix() 198 RS_PREFIX_SWEEP1(gl_LocalInvocationID.x) = h1_inc - h1_red; in rs_prefix() [all …]
|
H A D | scatter.glsl | 365 #define RS_PREFIX_LOAD(idx_) smem.extent[RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x + (idx… 366 #define RS_PREFIX_STORE(idx_) smem.extent[RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x + (idx… 397 const uint32_t smem_offset = RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x; 416 if (gl_LocalInvocationID.x < RS_RADIX_SIZE) 419 smem.extent[RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x] = 0; 683 const uint32_t hist_offset = gl_LocalInvocationID.x * 4; 720 const uint32_t smem_offset_h = RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x; 721 const uint32_t smem_offset_l = RS_SMEM_LOOKBACK_OFFSET + gl_LocalInvocationID.x; 766 if (gl_LocalInvocationID.x < RS_RADIX_SIZE) 770 const uint32_t red = smem.extent[RS_SMEM_HISTOGRAM_OFFSET + gl_LocalInvocationID.x]; [all …]
|
H A D | histogram.comp | 110 #define RS_HISTOGRAM_OFFSET(pass_) (RS_HISTOGRAM_BASE(pass_) + gl_LocalInvocationID.x * 4) 210 const uint32_t smem_offset = gl_LocalInvocationID.x; 227 if (gl_LocalInvocationID.x < RS_RADIX_SIZE) 230 smem.histogram[gl_LocalInvocationID.x] = 0; 258 const uint32_t smem_offset = gl_LocalInvocationID.x; 280 if (gl_LocalInvocationID.x < RS_RADIX_SIZE) 283 const uint32_t count = smem.histogram[gl_LocalInvocationID.x]; 339 umulExtended(gl_WorkGroupID.x * RS_BLOCK_KEYVALS + gl_LocalInvocationID.x,
|
/aosp_15_r20/external/deqp-deps/glslang/Test/baseResults/ |
D | web.comp.out | 9 …ryPoint GLCompute %main "main" %gl_NumWorkGroups %gl_WorkGroupID %gl_LocalInvocationID %gl_GlobalI… 24 OpName %gl_LocalInvocationID "gl_LocalInvocationID" 39 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId 90 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input 131 %78 = OpLoad %v3uint %gl_LocalInvocationID
|
D | spv.double.comp.out | 21 Name 33 "gl_LocalInvocationID" 31 Decorate 33(gl_LocalInvocationID) BuiltIn LocalInvocationId 57 33(gl_LocalInvocationID): 25(ptr) Variable Input 86 34: 24(ivec3) Load 33(gl_LocalInvocationID)
|
D | spv.ext.meshShaderTaskMem.mesh.out | 19 Name 11 "gl_LocalInvocationID" 32 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 48 11(gl_LocalInvocationID): 10(ptr) Variable Input 82 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
D | spv.meshShaderTaskMem.mesh.out | 19 Name 11 "gl_LocalInvocationID" 32 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 55 11(gl_LocalInvocationID): 10(ptr) Variable Input 89 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
D | spv.meshShaderSharedMem.mesh.out | 20 Name 11 "gl_LocalInvocationID" 29 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 45 11(gl_LocalInvocationID): 10(ptr) Variable Input 78 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
D | spv.310.comp.out | 32 Name 53 "gl_LocalInvocationID" 53 Decorate 53(gl_LocalInvocationID) BuiltIn LocalInvocationId 93 53(gl_LocalInvocationID): 52(ptr) Variable Input 124 55: 54(ptr) AccessChain 53(gl_LocalInvocationID) 33
|
/aosp_15_r20/external/angle/third_party/glslang/src/Test/baseResults/ |
H A D | web.comp.out | 9 …ryPoint GLCompute %main "main" %gl_NumWorkGroups %gl_WorkGroupID %gl_LocalInvocationID %gl_GlobalI… 24 OpName %gl_LocalInvocationID "gl_LocalInvocationID" 39 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId 90 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input 131 %78 = OpLoad %v3uint %gl_LocalInvocationID
|
H A D | spv.double.comp.out | 21 Name 33 "gl_LocalInvocationID" 31 Decorate 33(gl_LocalInvocationID) BuiltIn LocalInvocationId 57 33(gl_LocalInvocationID): 25(ptr) Variable Input 86 34: 24(ivec3) Load 33(gl_LocalInvocationID)
|
H A D | spv.ext.meshShaderTaskMem.mesh.out | 19 Name 11 "gl_LocalInvocationID" 32 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 48 11(gl_LocalInvocationID): 10(ptr) Variable Input 82 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
H A D | spv.meshShaderTaskMem.mesh.out | 19 Name 11 "gl_LocalInvocationID" 32 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 55 11(gl_LocalInvocationID): 10(ptr) Variable Input 89 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
H A D | spv.meshShaderSharedMem.mesh.out | 20 Name 11 "gl_LocalInvocationID" 29 Decorate 11(gl_LocalInvocationID) BuiltIn LocalInvocationId 45 11(gl_LocalInvocationID): 10(ptr) Variable Input 78 14: 13(ptr) AccessChain 11(gl_LocalInvocationID) 12
|
H A D | spv.310.comp.out | 32 Name 53 "gl_LocalInvocationID" 53 Decorate 53(gl_LocalInvocationID) BuiltIn LocalInvocationId 93 53(gl_LocalInvocationID): 52(ptr) Variable Input 124 55: 54(ptr) AccessChain 53(gl_LocalInvocationID) 33
|
/aosp_15_r20/external/executorch/backends/vulkan/tools/gpuinfo/glsl/ |
H A D | buf_bandwidth.glsl | 41 A[gl_LocalInvocationID[0]][0] = gl_LocalInvocationID[0]; 45 uint offset = (gl_WorkGroupID[0] * workgroup_width + gl_LocalInvocationID[0]) & addr_mask; 60 B[gl_LocalInvocationID[0]] = sum + zero;
|
/aosp_15_r20/external/deqp/external/vulkancts/data/vulkan/amber/rasterization/line_continuity/ |
H A D | line-strip.amber | 90 ivec2 p = ivec2(gl_LocalInvocationID) + ivec2(0, 128); 116 if (gl_LocalInvocationID.x < 3 && gl_LocalInvocationID.y < 3) 118 ivec2 p = pixel + ivec2(gl_LocalInvocationID) - ivec2(1);
|
H A D | polygon-mode-lines.amber | 92 ivec2 p = ivec2(gl_LocalInvocationID) + ivec2(0, 10); 118 if (gl_LocalInvocationID.x < 3 && gl_LocalInvocationID.y < 3) 120 ivec2 p = pixel + ivec2(gl_LocalInvocationID) - ivec2(1);
|
/aosp_15_r20/external/deqp/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/ |
H A D | workgroup_no_stride.amber | 31 OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID 37 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId 67 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input 80 %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
|
H A D | workgroup.amber | 33 OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID 39 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId 70 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input 83 %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
|
H A D | workgroup_bad_stride.amber | 32 OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID 38 OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId 69 %gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input 82 %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
|
/aosp_15_r20/external/deqp/external/vulkancts/data/vulkan/amber/compute/zero_initialize_workgroup_memory/ |
H A D | workgroup_size_8x4x4.amber | 32 uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y; 33 uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x; 34 uint idx_x = gl_LocalInvocationID.x;
|
H A D | workgroup_size_4x8x4.amber | 32 uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y; 33 uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x; 34 uint idx_x = gl_LocalInvocationID.x;
|
H A D | workgroup_size_4x4x8.amber | 32 uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y; 33 uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x; 34 uint idx_x = gl_LocalInvocationID.x;
|
H A D | workgroup_size_8x2x8.amber | 32 uint idx_z = gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y; 33 uint idx_y = gl_LocalInvocationID.y * gl_WorkGroupSize.x; 34 uint idx_x = gl_LocalInvocationID.x;
|