bbrezillon

arch=arm
autoload=no
baudrate=1500000
board=evb_rk3399
board_name=evb_rk3399
boot_a_script=load ${devtype} ${devnum}:${distro_bootpart} ${scriptaddr} ${prefix}${script}; source ${scriptaddr}
boot_efi_binary=if fdt addr ${fdt_addr_r}; then bootefi bootmgr ${fdt_addr_r};else bootefi bootmgr ${fdtcontroladdr};fi;load ${devtype} ${devnum}:${distro_bootpart} ${kernel_addr_r} efi/boot/bootaa64.efi; if fdt addr ${fdt_addr_r}; then bootefi ${kernel_addr_r} ${fdt_addr_r};else bootefi ${kernel_addr_r} ${fdtcontroladdr};fi
boot_extlinux=sysboot ${devtype} ${devnum}:${distro_bootpart} any ${scriptaddr} ${prefix}${boot_syslinux_conf}
boot_net_usb_start=usb start
boot_prefixes=/ /boot/

bbrezillon

nand-controller {
    /* 2 native IDs available */
    gpios-cs = <0>, <0>, <&pioA 2>;

    nand@0 {
        /* Native CS 0 */
        reg = <0>;
    };

    nand@2 {

bbrezillon

shader: MESA_SHADER_KERNEL
local-size: 12, 1, 1 (variable)
shared-size: 0
inputs: 2
outputs: 0
uniforms: 0
shared: 0
decl_var ubo INTERP_MODE_NONE uint[4] kernel_inputs (0, 0, 0)
decl_var ubo INTERP_MODE_NONE uint[11] kernel_work_properies (0, 0, 1)
decl_var shader_in INTERP_MODE_NONE uint64_t @0 (0.x, 0, 0)

bbrezillon

=====before vectorizer pass=====
shader: MESA_SHADER_COMPUTE
local-size: 0, 0, 0
shared-size: 0
inputs: 0
outputs: 0
uniforms: 0
shared: 0
decl_var shared INTERP_MODE_NONE uint[4] var
decl_function main (0 params)

bbrezillon

static nir_ssa_def *
build_round(nir_builder *nb, nir_ssa_def *src)
{
   nir_ssa_def *half = nir_imm_floatN_t(nb, 0.5, src->bit_size);
   nir_ssa_def *truncated = nir_ftrunc(nb, src);
   nir_ssa_def *remainder = nir_fsub(nb, src, truncated);

   return nir_bcsel(nb, nir_fge(nb, nir_fabs(nb, remainder), half),
                    nir_fadd(nb, truncated, nir_fsign(nb, src)), truncated);
}

bbrezillon

diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c
index 1e9b6d707e9..4c088bf6841 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -738,18 +738,7 @@ vtn_local_store(struct vtn_builder *b, struct vtn_ssa_value *src,
                 nir_deref_instr *dest, enum gl_access_qualifier access,
                 unsigned alignment)
 {
-   nir_deref_instr *dest_tail = get_deref_tail(dest);
-

bbrezillon

__kernel void test_fn(__local uint2 *sSharedStorage, __global uint2 *srcValues, __global uint *offsets, __global uint2 *destBuffer, uint alignmentOffset )    
{
    int tid = get_global_id( 0 );
 sSharedStorage[ offsets[tid] ] = (uint2)(uint)0;
 sSharedStorage[ offsets[tid] +1 ] =  sSharedStorage[ offsets[tid] ];
   barrier( CLK_LOCAL_MEM_FENCE );

    vstore2( srcValues[ tid ], offsets[ tid ], ( (__local uint *)sSharedStorage ) + alignmentOffset );

   barrier( CLK_LOCAL_MEM_FENCE );

bbrezillon

impl test_fn {
        block block_0:
        /* preds: */
        vec1 32 ssa_30 = load_const (0x00000001 /* 0.000000 */)
        vec2 32 ssa_24 = load_const (0x00000000 /* 0.000000 */, 0x00000000 /* 0.000000 */)
        vec1 64 ssa_15 = load_const (0x              20 /* 0.000000 */)
        vec1 64 ssa_12 = load_const (0x              20 /* 0.000000 */)
        vec1 64 ssa_0 = intrinsic load_param () (0) /* param_idx=0 */
        vec1 64 ssa_1 = deref_cast (uvec2 *)ssa_0 (shared uvec2)  /* ptr_stride=8 */
        vec1 64 ssa_2 = intrinsic load_param () (1) /* param_idx=1 */

bbrezillon

__kernel void test_fn(__local uint2 *sSharedStorage, __global uint2 *srcValues, __global uint *offsets, __global uint2 *destBuffer, uint alignmentOffset )    
{
    int tid = get_global_id( 0 );
 sSharedStorage[ offsets[tid] ] = (uint2)(uint)0;
 sSharedStorage[ offsets[tid] +1 ] =  sSharedStorage[ offsets[tid] ];
   barrier( CLK_LOCAL_MEM_FENCE );

    vstore2( srcValues[ tid ], offsets[ tid ], ( (__local uint *)sSharedStorage ) + alignmentOffset );

   barrier( CLK_LOCAL_MEM_FENCE );

bbrezillon

shader: MESA_SHADER_KERNEL
local-size: 1, 1, 1 (variable)
shared-size: 0
inputs: 5
outputs: 0
uniforms: 0
shared: 0
decl_var uniform INTERP_MODE_NONE readonly isampler1D @0 (0, 0, 0)
decl_var uniform INTERP_MODE_NONE readonly usampler1D @1 (0, 0, 1)
decl_var uniform INTERP_MODE_NONE sampler @2 (0, 0, 0)