text:
numerics issue with bert on cpu

What happened?
When running bert-based-uncased we seemed to be hitting a numerics issue:
https://github.com/nod-ai/SHARK/blob/main/tank/tf/hf_masked_lm/bert-base-uncased_test.py
IREE_SAVE_TEMPS  is at iree-shared-files/nod-perf/anush/iree_save_temps_bert_cpu
https://storage.googleapis.com/iree-shared-files/nod-perf/anush/iree_save_temps_bert_cpu/core-input.mlir
https://storage.googleapis.com/iree-shared-files/nod-perf/anush/iree_save_temps_bert_cpu/core-command-line.txt
E       assert True == False                                                                                E        +  where False = compare_tensors_tf(<tf.Tensor: shape=(1, 16, 30522), dtype=float32, numpy=\narray([[[ -7.290994 ,  -7.257680
4,  -7.2415266, ...,  -6.464171 ,\n          -6.3598876,  -4.309055 ],\n        [ -8.907964 ,  -8.83476  ,  -8.829174 , ...,  -8.76890
6 ,\n          -8.046408 ,  -6.274063 ],\n        [-14.354682 , -14.289836 , -14.320234 , ..., -13.399221 ,\n         -10.088572 , -11
.495969 ],\n        ...,\n        [ -8.209971 ,  -8.197388 ,  -8.366645 , ...,  -7.885937 ,\n          -8.4152775,  -3.0067217],\n        [ -7.3771143,  -7.7098002,  -7.6014137, ...,  -8.009481 ,\n          -7.0031247,  -5.9178457],\n        [ -5.960579 ,  -6.303993 ,  -6.0986414, ...,  -6.343738 ,\n          -5.770889 ,  -4.7784796]]], dtype=float32)>, array([[[ -7.291651 ,  -7.2583814,  -7.2421713
, ...,  -6.464848 ,\n          -6.3605504,  -4.309505 ],\n        [ -8.909623 ,  -8.8363695,  -8.830784 , ...,  -8.770344 ,\n
 -8.047972 ,  -6.2748594],\n        [-14.355753 , -14.291211 , -14.321319 , ..., -13.400625 ,\n         -10.089375 , -11.49576  ],\n
      ...,\n        [ -8.208423 ,  -8.195899 ,  -8.3648815, ...,  -7.8849854,\n          -8.413467 ,  -3.008876 ],\n        [ -7.37734
75,  -7.7098885,  -7.601365 , ...,  -8.009851 ,\n          -7.0026016,  -5.9188967],\n        [ -5.9612074,  -6.304814 ,  -6.09935  ,
...,  -6.3439813,\n          -5.772274 ,  -4.775791 ]]], dtype=float32))

Steps to reproduce your issue

Go to '...'
Click on '....'
Scroll down to '....'
See error

What component(s) does this issue relate to?
No response
Version information
No response
Additional context
No response

I've gotten some updated reproduction resources here. This is for numerics not matching between IREE and TF results running inference on bert-base-uncased (static) on a VM with a2-highgpu-1g machine type (intel cascade lake). It's imporant to note that the same issue is not reproducible (the IREE and tf results match within tolerance) on N1 (intel haswell) and N2 (intel cascade lake) machine types, which suggests that the CPU platform shouldn't be an issue.
Both of the following cases fail to validate (np.allclose returns false between IREE result and TF result with rtol=1e-02 and atol=1e-03):
Without the iree-llvm-target-cpu-features=host flag:
IREE_SAVE_TEMPS is at gs://iree-shared-files/nod-perf/ean/bert_repro_nohost/iree_temps/
Result tensor comparison, function inputs, input MLIR, and .vmfb are at gs://iree-shared-files/nod-perf/ean/bert_repro_nohost/
With the iree-llvm-target-cpu-features=host flag:
IREE_SAVE_TEMPS is at gs://iree-shared-files/nod-perf/ean/bert_repro/iree_temps/
Result tensor comparison, function inputs, input MLIR, and .vmfb are at gs://iree-shared-files/nod-perf/ean/bert_repro/
I was able to reproduce the IREE results though IREE CLI via the following command:
--module_file=./bert_repro/tensorflow_forward_cpu.vmfb \
--device=local-task \
--entry_function=forward \
--function_input=@./bert_repro/input1.npy \
--function_input=@./bert_repro/input2.npy```

Thanks @monorimet . Ill take a look at it in the afternoon

A few more notes from our discussion on Discord:

The expected results generated through TensorFlow are different between machine configurations (enough to stretch outside of allowed tolerance, 1e-02 ~ 1e-4), likely due to OneDNN optimizations. This could indicate that the IREE results are correct but I don't know if we can say for sure.

Without the iree-llvm-target-cpu-features=host flag, IREE results on all configurations are the same. This suggests that the IREE results from the a2-highgpu-1g + NVIDIA A100 configuration without the `target_cpu_features=host' flag would match the expected results from TF on the n1/Haswell machine config.

We are experiencing similar numerics issues with Vulkan. (could be separate issue, filing separate bugs just in case)

We are experiencing the same issue with other bert-based language models (albert-base-v2, camembert-base, etc.)

I'd like to see if the IREE result from our a2 setup validates with the TF result from the n1/haswell setup. That would further indicate that there's a big results diff coming from the TensorFlow opts. I'll do some more looking tonight.

#9524 might be related actually. There is a first-level work around that will attempted based on the vulkan bug (https://github.com/google/iree/issues/9524) . If that fixes CPU and CUDA as well, then the issue might be related to memref.subview handling during lowering to LLVM (which shows up as a compile time error on Vulkan side).

hmm.. when the vulkan issue was filed I think CPU was fine but don't remember now.

just going off of this We are experiencing the same issue with other bert-based language models (albert-base-v2, camembert-base, etc.)

Ill look into the CPU side, but in the meantime can you post some info on the difference for the CUDA side...

@hanhanW as well

I'm able to narrow it down to dispatch_24. It looks like an accumulation of noise from f32 optimization.
#map0 = affine_map<(d0, d1) -> (d0, d1)>
#map3 = affine_map<(d0, d1) -> (d0)>
func.func @forward_dispatch_24(%0: tensor<16x768xf32>, %1: tensor<16xf32>) -> tensor<16xf32> {
  %cst = arith.constant -0.000000e+00 : f32
  %2 = linalg.init_tensor [16] : tensor<16xf32>
  %3 = linalg.fill ins(%cst : f32) outs(%2 : tensor<16xf32>) -> tensor<16xf32>
  %4 = linalg.generic {indexing_maps = [#map0, #map3, #map3], iterator_types = ["parallel", "reduction"]}
    ins(%0, %1 : tensor<16x768xf32>, tensor<16xf32>)
    outs(%3 : tensor<16xf32>) {
  ^bb0(%arg3: f32, %arg4: f32, %arg5: f32):
    %5 = arith.subf %arg3, %arg4 : f32
    %6 = arith.mulf %5, %5 : f32
    %7 = arith.addf %6, %arg5 : f32
    linalg.yield %7 : f32
  } -> tensor<16xf32>
  return %4 : tensor<16xf32>
}

dispatch_24 is called several times. The input of dispatch_24 are quite different between two setup. The max diff of input1 is increasing from 1e-6 to 1e-4, see log.
Here is the dump in flagfile format from logs.
If we compile the dispatch_24 w/o host features, we'll get
EXEC @forward_dispatch_24
result[0]: hal.buffer_view
16xf32=
2056.357422 996.230835 1129.984375 847.100098
1054.056396 1109.782959 834.492310 891.951355
2034.976196 81371.328125 1890.395630 1775.567505
2241.176758 2055.166504 2253.345215 2316.949951

If we compile the dispatch_24 w/ host features, we'll get
EXEC @forward_dispatch_24
result[0]: hal.buffer_view
16xf32=
2056.357178 996.230896 1129.984619 847.100342
1054.056641 1109.782593 834.492310 891.951660
2034.975952 81371.320312 1890.394775 1775.567139
2241.177734 2055.167236 2253.344482 2316.949219

Look into the 10-th number (i.e., the third number in third row), the diff is ~0.008  (81371.328125 versus 81371.320312).

The albert_base_v2_temps case looks fine to me. I checked the outputs between w/ and w/o host features using rtol=1e-02, atol=1e-03 config, and it passes. I'd claim that the diff is reasonable in this case. Do we want someone who is more familiar with models to verify if they are correct?
I'm able to check other models, but it takes some time. It would be good if someone can help verify the noise is acceptable or not.

awesome. if you have a pip package with a fix we can install it in the shark.venv and just run all tests. We can also adjust tolerance if required.

wait wait :D
We saw the same on N1 and non a2-highgpu-1g instances. It is only for the a2-highgpu instance we saw the difference. So lets test on that ?

so all 6 models are within tolerance ?

No, I haven't tried on all the models. Since we have small reproducer, I'd like to put more focus on the smaller repro. Will give it a shot after we understand why.
===
I tested with several configurations and found that the root cause is tiling sizes. The host features affect vector size, which introduces different tiling sizes. E.g., one is [8 ,4, 4], and the other is [16, 16, 16].
Different tiling sizes produce different output.
If I set the size to 1 (which is scalar implementation), I'll get81371.257812.
If I set the size to 4, I'll get 81371.328125.
If I set the size to 16, I'll get 81371.320312
I looked into the asm, and found that the use of single-precision and double-precision is mixed. It results in numerical issue. E.g.,
This is the graph for tile_sizes = [[8, 0], [1, 0], [0, 1]] config: https://onlinedisassembler.com/odaweb/9pi4SjPV/0

This is the graph for tile_sizes = [[8, 0], [4, 0], [0, 4]] config: https://onlinedisassembler.com/odaweb/WlI1co65

The scalar version uses xmm register to store single values, and uses subss instruction for computation. The vectorized version uses xmm register to store packed values. and uses subps instruction for computation.
subss performs subtraction on single-precision floating-point values, and store the double-precision floating-point result in the low doubleword of the destination operand.
subps performs a SIMD subtract of the packed single-precision floating-point values and store the packed single-precision floating-point values to destination.
One results in double-precision values and the other results in packed single-precision values. This might introduces numerical issues. The same thing happens between tile_size=4 and tile_size=16 cases. Thus, they are all correct to me.

Really nice analysis Hanhan! Where is the double precision being introduced? Is it in MLIR or when we lower to LLVM something in LLVM compilation is promoting to 64 bit?

Tracked down the issue. The vectorized path for reductions (and I suspect gemm as well) goes through the following intrinsic
 %131 = "llvm.intr.vector.reduce.fadd"(%1, %130) {reassoc = false} : (f32, vector<16xf32>) -> f32

which lowers to
%423 = call float @llvm.vector.reduce.fadd.v16f32(float 0.000000e+00, <16 x float> %422), !dbg !13

According to LangRef since this is lowering to the reassoc = false mode the order of operations is different. That explains the difference in values.
Ideally we would have a flag to switch reassoc=false to reassoc=true, requires more surgery though upstream. So now the difference in values makes sense. The answer is to increase the tolerance...
@powderluv closing this issue itself, since this is expected behavior.

I'm a bit confused why we are closing this -- it sounds like there is work to be done to manage the reassoc attribute? Not saying your wrong but I think I'll need some convincing on why tile size changes are producing such different code paths with respect to precision. I know fp optimizations are prone to cause numerical differences in extreme cases, but this doesn't seem that.

The different tile sizes produce different vector sizes. So then it's either a 16-way reduce or an 8-way reduce or 4-way reduce.
If we drop the reassoc attribute then the reduction effectively gets scalarized. We probably want to keep the vectorization and increase tolerance.
I closed it cause I don't see a correctness issue per se. We can try to have a mode that doesn't have a difference, but that is as good as a flag that just disables vectorization. Did you have something else in mind here?

I think I can add an option to mark the attribute to true, and see if it's still an issue.

I think we want an option to do that -- so we have a baseline correct solution or if you are ok we can go fast fast but can't guarantee the outcomes / tolerances and it left to the user to verify it is ok.
An end customer is going to compare us to TF / PT models and we are going to be dealing with customers where this tolerance makes $$ differences.
Also similar discussions on why there are options to turn of non-determinism in PyTorch etc.
https://discuss.pytorch.org/t/what-is-the-differenc-between-cudnn-deterministic-and-cudnn-benchmark/38054 and https://pytorch.org/docs/stable/notes/randomness.html
This will also be a factor in the tuner - a customer may want to tune for better performance but not at the cost of lower accuracy (because we chose a different tile size) and another may be ok with the variance.

The option is to not use iree-llvm-target-cpu-features=host. That is effectively what setting the reassociation to true (which conversely implies don't use reassociation) would do, scalarizes the code. Worth trying to make sure this is the root cause, but I don't see a need for an additional flag.

I see. Using iree-llvm-target-cpu-features=host means that we allow the compiler being aggressive on performance. In this case, the reassoc attribute is managed by the flag.
I'm still confused why {reassoc = false} is set even without host features. Maybe we're missing some controls on it. I'll take a look at it.

It is always set. That seems to be the upstream lowering choice. The difference is vector length. It's either 4 length or 16 length. Higher the length more the change in value. I'll confirm in the morning (if you haven't gotten to it already) that setting it to true fixes the issue.

Okay, it's set to false by default.
void populateVectorToLLVMConversionPatterns(
    LLVMTypeConverter &converter, RewritePatternSet &patterns,
    bool reassociateFPReductions = false, bool force32BitVectorIndices = false);

I just tried to set it true, but the outputs are identical. It does not fix the issue. Also, the result is different from scalar implementation (i.e., without vectorization). I'll dump the LLVM module, and check if the attribute is propagated well..

I tried increasing the tolerance for validating inference through IREE vs. through Tensorflow. The smallest tolerance for which the tests will pass is rtol = 1e-02, atol = 1e-01. I doubt that such a large tolerance is sufficient to validate inference results for these models.
We were previously validating with rtol = 1e-02, atol = 1e-03 (tests fail with these values).
For reference, HuggingFace model tests validate with rtol = default and atol = 1e-04.
The default atol and rtol is 10 * eps, where eps is the smallest representable positive number such that 1 + eps != 1. This is about 1.2e-6 in 32bit, 2.22e-15 in 64bit, and 0.00977 in 16bit
MLPerf compares the accuracies of inference results, but expects prediction accuracies to be 99% similar.

Here is how the same model is compared in HuggingFace: https://github.com/huggingface/transformers/blob/main/tests/models/albert/test_modeling_tf_albert.py#L318
tf.debugging.assert_near(output[:, :3, :3], expected_slice, atol=1e-4)

The same issue occurs without the iree-llvm-target-cpu-features=host flag.
It also occurs with the dynamic cases for the models listed with repros above on CPU.
For example, I have another set of repro artifacts here for albert dynamic without the ...cpu-features=host flag:
gs://iree-shared-files/nod-perf/ean/albert_nohost_temps/

Also albert dynamic with CUDA without the flag:
gs://iree-shared-files/nod-perf/ean/albert_nohost_cuda_temps/

I'm working on diagnosing our CUDA errors as well -- I expect all of the models having this issue on CPU to have the same issue on CUDA, but will verify soon.

As reported above, we're able to reproduce the issue without iree-llvm-target-cpu-features=host. The root cause it that different tiling sizes are producing different outputs. We're still investigating why.
@MaheshRavishankar I think the default behavior is what we want. It will produce a scalarized reduction if reassoc = false. If reassoc = false, the intrinsic call does not have the ‘reassoc’ flag set.
See the examples:
%unord = call reassoc float @llvm.vector.reduce.fmul.v4f32(float 1.0, <4 x float> %input) ; relaxed reduction
%ord = call float @llvm.vector.reduce.fmul.v4f32(float %start_value, <4 x float> %input) ; sequential reduction

We want ordered computation. And the IR we generated looks good to me.
%423 = call float @llvm.vector.reduce.fadd.v16f32(float 0.000000e+00, <16 x float> %422), !dbg !13

Here is my guess -- the tiling on reduction dim affects the order of computation.
If we run the computation sequentially. We're doing
int sum = 0;
for (int i = 0; i < n; ++i) {
  sum += in[i];
}

If we tile the reduction dim, the order changes. We'll apply reduction on each tile and sum them up.
int sum = 0;
for (int i = 0; i < n/4; ++i) {
  int inner_sum = 0;
  for (int j = 0; j < 4; ++j) {
    inner_sum += in[j];
  }
  sum += inner_sum;
}

The below IR at LLVM level shows that it applies reduction on 4 floating-point values, and adds them to destination.
  %127 = insertelement <4 x float> %122, float %126, i64 2, !dbg !13
  %128 = extractelement <4 x float> %90, i64 3, !dbg !13
  %129 = fmul <4 x float> %112, %112, !dbg !13
  %130 = tail call float @llvm.vector.reduce.fadd.v4f32(float 0.000000e+00, <4 x float> %129), !dbg !13
  %131 = fadd float %128, %130, !dbg !13
  %132 = insertelement <4 x float> %127, float %131, i64 3, !dbg !13

Does it explain why the outputs are different when we have different tiling sizes? If so, we'll have to figure out a fix or revisit our codegen pipeline..

Tiling on reduction dimension does not change the order of operations AFAIK. If we don't tile the reduction dimension it won't get vectorized along the reduction dimension
@hanhanW did you check the LLVM IR to see if the reassoc flag is propagated

Looked at the LangRef again, my reading was off. I'll take a look again when I get online

I am pretty sure that I can come up with values for in that produce such different value even in the two C versions (tiled and not tiled) because in fp you get more precision around zero.
I believe there are some practical ways to stabilize such partial sums but it has been a while for me and will need to do some research.

Tiling shouldn't be using init value as zero for every tile. It should use the result of the previous tile as the unit value for the next tile. That happens in Linalg. If the vector instructions are instead using zero as the initialization value that would be wrong indeed

I think that if you want to be accumulating partial sums like that, you need to be accumulating at higher precision or using a more expensive algorithm. But as Mahesh notes, of not parallelizing the reduction, accumulating in place is the cheapest.
(We have actual fp experts in the team who should weigh in if more nuanced than these options)

Following up now based on some comments above. I might have found the root cause. Here is the IR before the vectorization
    %8 = scf.for %arg1 = %c0 to %c8 step %c4 iter_args(%arg2 = %5) -> (tensor<8xf32>) {
      %9 = tensor.extract_slice %7[%arg1] [4] [1] : tensor<8xf32> to tensor<4xf32>
      %10 = tensor.extract_slice %arg2[%arg1] [4] [1] : tensor<8xf32> to tensor<4xf32>
      %11 = linalg.fill {__internal_linalg_transform__ = "1"} ins(%cst : f32) outs(%10 : tensor<4xf32>) -> tensor<4xf32>
      %12 = scf.for %arg3 = %c0 to %c768 step %c4 iter_args(%arg4 = %11) -> (tensor<4xf32>) {
        %14 = tensor.extract_slice %6[%arg1, %arg3] [4, 4] [1, 1] : tensor<8x768xf32> to tensor<4x4xf32>
        %15 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"]} ins(%14, %9 : tensor<4x4xf32>, tensor<4xf32>) outs(%arg4 : tensor<4xf32>) attrs =  {__internal_linalg_transform__ = "1", lowering_config = #iree_codegen.lowering_config<tile_sizes = [[8, 0], [4, 0], [0, 4]]>} {
        ^bb0(%arg5: f32, %arg6: f32, %arg7: f32):
          %16 = arith.subf %arg5, %arg6 : f32
          %17 = arith.mulf %16, %16 : f32
          %18 = arith.addf %17, %arg7 : f32
          linalg.yield %18 : f32
        } -> tensor<4xf32>
        scf.yield %15 : tensor<4xf32>
      }
      %13 = tensor.insert_slice %12 into %arg2[%arg1] [4] [1] : tensor<4xf32> into tensor<8xf32>
      scf.yield %13 : tensor<8xf32>
    }

Things to notice, the %arg4 is used as the outs operand to the linalg.generic. So after one iteration of the inner loop, the result of the linalg.generic is yielded and fed back into the loop. So the order of computation I see for the reduction is
sum = init_value + val[0] + val[1] + val[2] .....

and after vectorization
    %8 = scf.for %arg1 = %c0 to %c8 step %c4 iter_args(%arg2 = %5) -> (tensor<8xf32>) {
      %9 = tensor.extract_slice %arg2[%arg1] [4] [1] : tensor<8xf32> to tensor<4xf32>
      %10 = vector.transfer_write %cst, %9[%c0] {in_bounds = [true]} : vector<4xf32>, tensor<4xf32>
      %11 = scf.for %arg3 = %c0 to %c768 step %c4 iter_args(%arg4 = %10) -> (tensor<4xf32>) {
        %13 = vector.transfer_read %6[%arg1, %arg3], %cst_0 {in_bounds = [true, true]} : tensor<8x768xf32>, vector<4x4xf32>
        %14 = vector.transfer_read %7[%arg1], %cst_0 {in_bounds = [true]} : tensor<8xf32>, vector<4xf32>
        %15 = vector.broadcast %14 : vector<4xf32> to vector<4x4xf32>
        %16 = vector.transpose %15, [1, 0] : vector<4x4xf32> to vector<4x4xf32>
        %17 = vector.transfer_read %arg4[%c0], %cst_0 {in_bounds = [true]} : tensor<4xf32>, vector<4xf32>
        %18 = arith.subf %13, %16 : vector<4x4xf32>
        %19 = vector.contract {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"], kind = #vector.kind<add>} %18, %18, %cst_1 : vector<4x4xf32>, vector<4x4xf32> into vector<4xf32>
        %20 = arith.addf %19, %17 : vector<4xf32>
        %21 = vector.transfer_write %20, %arg4[%c0] {in_bounds = [true]} : vector<4xf32>, tensor<4xf32>
        scf.yield %21 : tensor<4xf32>
      }
      %12 = tensor.insert_slice %11 into %arg2[%arg1] [4] [1] : tensor<4xf32> into tensor<8xf32>
      scf.yield %12 : tensor<8xf32>
    }

AFAICS there is a slight change in ordering here. The vector.contract starts from an initial value of zero for every iteration. So the order of computation I see for the reduction is
sum = (0 + val[0] + val[1] + val[2] + val[3]) + init_value + (0 + val[4] + val[5] + val[6] + val[7]) + ...

This is associatively different. @nicolasvasilache and @ThomasRaoux if this makes sense to you. Dont know how easy the fix is.

Actually it is a little bit more convoluted than that because of arith.addf %19, %17.  It is
sum = (0 + val[n-4] + val[n-3] + val[n-2] + val[n-1]) + (0 + val[n-8] + val[n-7] + val[n-6] + val[n-5]) + .... + (val[3] + val[2] + val[1] + val[0]) + init_value.

so you would expect addf and contract order to be swapped?

That would do it. The multiple accumulation into zero is one issue and the different associativity is another. They are probably separable from a numerics perspective.

If I am reading this correctly, I would expect the arith.addf to not be there. I'd expect
    %8 = scf.for %arg1 = %c0 to %c8 step %c4 iter_args(%arg2 = %5) -> (tensor<8xf32>) {
      %9 = tensor.extract_slice %arg2[%arg1] [4] [1] : tensor<8xf32> to tensor<4xf32>
      %10 = vector.transfer_write %cst, %9[%c0] {in_bounds = [true]} : vector<4xf32>, tensor<4xf32>
      %11 = scf.for %arg3 = %c0 to %c768 step %c4 iter_args(%arg4 = %10) -> (tensor<4xf32>) {
        %13 = vector.transfer_read %6[%arg1, %arg3], %cst_0 {in_bounds = [true, true]} : tensor<8x768xf32>, vector<4x4xf32>
        %14 = vector.transfer_read %7[%arg1], %cst_0 {in_bounds = [true]} : tensor<8xf32>, vector<4xf32>
        %15 = vector.broadcast %14 : vector<4xf32> to vector<4x4xf32>
        %16 = vector.transpose %15, [1, 0] : vector<4x4xf32> to vector<4x4xf32>
        %17 = vector.transfer_read %arg4[%c0], %cst_0 {in_bounds = [true]} : tensor<4xf32>, vector<4xf32>
        %18 = arith.subf %13, %16 : vector<4x4xf32>
        %19 = vector.contract {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"], kind = #vector.kind<add>} %18, %18, %17 : vector<4x4xf32>, vector<4x4xf32> into vector<4xf32>
        %21 = vector.transfer_write %20, %arg4[%c0] {in_bounds = [true]} : vector<4xf32>, tensor<4xf32>
        scf.yield %21 : tensor<4xf32>
      }
      %12 = tensor.insert_slice %11 into %arg2[%arg1] [4] [1] : tensor<4xf32> into tensor<8xf32>
      scf.yield %12 : tensor<8xf32>
    }

right, I agree this is wrong. I'm surprised the contract and add are not merged back together though as we usually we run CanonicalizeContractAdd pattern right after which will merge those back. Is this is not happening in the CPU pass?
Not that it makes it right but I would assume those two would cancel each other and make the IR what you expect.

Here is the dump after all https://gist.github.com/MaheshRavishankar/bd2dc0340a0cf50ffa71d5402a2c145f . We probably dont have that pattern explicitly. CPU side is just using the CodegenStrategy upstream. So not sure how all those pieces are put together there.

@ThomasRaoux if its not something that would just take a few mins, I can try to take a stab at fixing it.

I think the merging back has the same issues w.r.t changing order of operations. I think it might be better to fix the vectorization of linalg.generic. Not familiar with that code so dont know how easy that is.

Yes as I mentioned it doesn't make the transformation right but I would expect the IR to be what you expect and therefore wouldn't cause a precision difference in this case.

This is a pretty important invariant to maintain -- without knowing that code and making a context free assessment, I would expect the transform itself to be correct and not rely on an after the fact cleanup.

yes this is definitely a bug, I'm just saying I'm not sure this is the root cause of the precision difference.
I do see the ops being merged in the dump:
scf.for %arg0 = %3 to %c16 step %4 {
    %5 = flow.dispatch.tensor.load %2, offsets = [%arg0], sizes = [8], strides = [1] : !flow.dispatch.tensor<writeonly:16xf32> -> tensor<8xf32>
    %6 = flow.dispatch.tensor.load %0, offsets = [%arg0, 0], sizes = [8, 768], strides = [1, 1] : !flow.dispatch.tensor<readonly:16x768xf32> -> tensor<8x768xf32>
    %7 = flow.dispatch.tensor.load %1, offsets = [%arg0], sizes = [8], strides = [1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<8xf32>
    %8 = scf.for %arg1 = %c0 to %c8 step %c4 iter_args(%arg2 = %5) -> (tensor<8xf32>) {
      %9 = vector.transfer_read %7[%arg1], %cst_0 {in_bounds = [true]} : tensor<8xf32>, vector<4xf32>
      %10 = vector.broadcast %9 : vector<4xf32> to vector<4x4xf32>
      %11 = vector.transpose %10, [1, 0] : vector<4x4xf32> to vector<4x4xf32>
      %12 = scf.for %arg3 = %c0 to %c768 step %c4 iter_args(%arg4 = %cst) -> (vector<4xf32>) {
        %14 = vector.transfer_read %6[%arg1, %arg3], %cst_0 {in_bounds = [true, true]} : tensor<8x768xf32>, vector<4x4xf32>
        %15 = arith.subf %14, %11 : vector<4x4xf32>
        %16 = vector.contract {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"], kind = #vector.kind<add>} %15, %15, %arg4 : vector<4x4xf32>, vector<4x4xf32> into vector<4xf32>
        scf.yield %16 : vector<4xf32>
      }
      %13 = vector.transfer_write %12, %arg2[%arg1] {in_bounds = [true]} : vector<4xf32>, tensor<8xf32>
      scf.yield %13 : tensor<8xf32>
    }
    ```

> @ThomasRaoux if its not something that would just take a few mins, I can try to take a stab at fixing it.

The fix is not trivial, either way is fine with me, feel free to go for it or pass it on to me.

yeah, that seems to be right. Ill keep looking. Maybe found an issue that is not the cause here. Would be still good to fix it. If you think the fix is not trivial I'd very much appreciate if you could do it (Cause it will be more non-trivial for me). Might not help this case though. Let me see.

yeah, that seems to be right. Ill keep looking. Maybe found an issue that is not the cause here. Would be still good to fix it. If you think the fix is not trivial I'd very much appreciate if you could do it (Cause it will be more non-trivial for me). Might not help this case though. Let me see.

Definitely I can work on a fix. Let me know what you find and I'll prioritize based on that.

Next candidate
Before lower vectors pass
    %10 = scf.for %arg1 = %c0 to %c768 step %c4 iter_args(%arg2 = %cst) -> (vector<4xf32>) {
      %11 = vector.transfer_read %5[%arg0, %arg1], %cst_0 {in_bounds = [true, true]} : memref<8x768xf32, affine_map<(d0, d1)[s0] -> (d0 * 768 + s0 + d1)>>, vector<4x4xf32>
      %12 = arith.subf %11, %9 : vector<4x4xf32>
      %13 = vector.contract {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"], kind = #vector.kind<add>} %12, %12, %arg2 : vector<4x4xf32>, vector<4x4xf32> into vector<4xf32>
      scf.yield %13 : vector<4xf32>
    }

The ordering here is
sum = init + a[0] * b[0] + a[1] * b[1] + a[2] * b[2] ....

After lower vectors pass
    %10 = scf.for %arg1 = %c0 to %c768 step %c4 iter_args(%arg2 = %cst_0) -> (vector<4xf32>) {
      %11 = vector.transfer_read %5[%arg0, %arg1], %cst_1 {in_bounds = [true, true]} : memref<8x768xf32, affine_map<(d0, d1)[s0] -> (d0 * 768 + s0 + d1)>>, vector<4x4xf32>
      %12 = arith.subf %11, %9 : vector<4x4xf32>
      %13 = vector.extract %12[0] : vector<4x4xf32>
      %14 = vector.extract %12[0] : vector<4x4xf32>
      %15 = vector.extract %arg2[0] : vector<4xf32>
      %16 = arith.mulf %13, %14 : vector<4xf32>
      %17 = vector.reduction <add>, %16 : vector<4xf32> into f32
      %18 = arith.addf %17, %15 : f32
      %19 = vector.insert %18, %cst [0] : f32 into vector<4xf32>
      %20 = vector.extract %12[1] : vector<4x4xf32>
      %21 = vector.extract %12[1] : vector<4x4xf32>
      %22 = vector.extract %arg2[1] : vector<4xf32>
      %23 = arith.mulf %20, %21 : vector<4xf32>
      %24 = vector.reduction <add>, %23 : vector<4xf32> into f32
      %25 = arith.addf %24, %22 : f32
      %26 = vector.insert %25, %19 [1] : f32 into vector<4xf32>
      %27 = vector.extract %12[2] : vector<4x4xf32>
      %28 = vector.extract %12[2] : vector<4x4xf32>
      %29 = vector.extract %arg2[2] : vector<4xf32>
      %30 = arith.mulf %27, %28 : vector<4xf32>
      %31 = vector.reduction <add>, %30 : vector<4xf32> into f32
      %32 = arith.addf %31, %29 : f32
      %33 = vector.insert %32, %26 [2] : f32 into vector<4xf32>
      %34 = vector.extract %12[3] : vector<4x4xf32>
      %35 = vector.extract %12[3] : vector<4x4xf32>
      %36 = vector.extract %arg2[3] : vector<4xf32>
      %37 = arith.mulf %34, %35 : vector<4xf32>
      %38 = vector.reduction <add>, %37 : vector<4xf32> into f32
      %39 = arith.addf %38, %36 : f32
      %40 = vector.insert %39, %33 [3] : f32 into vector<4xf32>
      scf.yield %40 : vector<4xf32>
    }

I think this goes from
sum = (a[n-4] * b[n-4] + a[n-3] * b[n-3] + a[n-2] * b[n-2] + a[n-1] * b[n-1]) + (a[n-8] * b[n-8] + a[n-7] * b[n-7] + a[n-6] * b[n-6] + a[n-5] * b[n-5]) + ..... + (a[0] * b[0] + a[1] * b[1] + a[2] * b[2] + a[3] * b[3]) + init.

This is because of arith.addf %17, %15 (and so on). I think those need to be flipped to arith.addf %15, %17
@ThomasRaoux does that compute (I am just reading the IR so I might be off as well)

AFAICS, this survives all the way.

Interesting, it does re-order instructions but I don't think there is another way to use the horizontal reduction. This is what we discussed earlier. If we do horizontal reduction I believe we end up doing partial sums, I don't see a different way to do it.

I think just swapping the order of operands in arith.addf should fix it.

Ill try to verify that indeed that fixes it.

you mean changing arith.addf %17, %15 -> to arith.addf %15, %17? This wouldn't change anything as addf is commutative

IIUC this ordering is changing the associativity as well. If you flip the order of operands the order would be what you expect I think
sum = init + a[0] * b[0] + a[1] * b[1] + .....

hmm Associativity and commutativity are independent:
sum = init + (a[0] * b[0] + a[1] * b[1] + .....) is the same as:
sum = (a[0] * b[0] + a[1] * b[1] + .....) + init
but different than:
sum =  ((init + a[0] * b[0]) + a[1] * b[1]) + .....

are you saying
sum = (a[n-4] * b[n-4] + a[n-3] * b[n-3] + a[n-2] * b[n-2] + a[n-1] * b[n-1]) + (a[n-8] * b[n-8] + a[n-7] * b[n-7] + a[n-6] * b[n-6] + a[n-5] * b[n-5]) + ..... + (a[0] * b[0] + a[1] * b[1] + a[2] * b[2] + a[3] * b[3]) + init.

is same as
sum = init + a[0] * b[0] + a[1] * b[1] + .....

well depends in what order the 3 operations in your first snippet (and it doesn't depend on the order of operands in individual addf).
What is the original semantic is:
((init + a[0]*b[0]) + a[1] * b[1]) + ...
What we generate is:
init + (a[0]*b[0]) + a[1] * b[1] + ...)
which are different due to the order of operations.

I'm a bit confused, probably we should move the discussion to discord or gvc

To summarize what I meant earlier, I think if we want to do horizontal reduction we don't have an alternative and it will change precision, if we don't want that then we can just serialize the reduction. It would be interesting to see the precision difference this cause.
If it is what is causing the difference we are looking for then the problem is really that this is an unsafe floating point optimization

Lunchtime musing: if we do end up with something the user needs to control to get a major performance advantage (and parallelizing horizontal reductions sounds like one, if we can't otherwise algorithmically avoid them) an option would be to push it up the stack - "use jax.foo_with_reduced_precision if you have a performance sensitive operation" -> mhlo.foo {reduced_precision} -> etc - then we can guarantee the tradeoff is consistently applied and in cases of just "mhlo.foo" always serialize/be safe. Given that XLA and others make these decisions seemingly at random this would at least surface the control to the user/framework so they can make the decision locally vs. a compiler flag that would change all occurrences in the entire input program. We could still provide a flag (fast-math style) that blanket applied the attribute but wouldn't enforce it.
That is to say if we have to choose between performance and precision we could instead choose to make the user choose ;P

I agree, if we go with flags we would want per operation flag. This is was has been done on llvm for instance. This would allow the front end to give different kind of granularity to user based on need. It could be a compiler setting it for all ops or it could be controlled fine grained by user.
There has been discussions in mlir discourse about those kind of flags but it will require heavy plumbing across the stack.

Hopefully this is the real root cause (third time the charm)
    %60 = vector.reduction <add>, %59 : vector<4xf32> into f32
    %61 = arith.addf %60, %58 : f32

%59 is the "input" %58 is the accumulator. So this seems OK because of commutativity.
This gets lowered to LLVM Dialect this way
    %302 = llvm.mlir.constant(0.000000e+00 : f32) : f32
    %303 = "llvm.intr.vector.reduce.fadd"(%302, %301) {reassoc = false} : (f32, vector<4xf32>) -> f32
    %304 = llvm.fadd %303, %300  : f32

%59 -> %301 and %58 -> %300 . So again this doing reduction into zero initial value and then adding the accumulation. That changes the precision.
I think the fix is vector.reduction <add> needs to take the init value as well as argument. Then we can again avoid doing the arith.addf in the first snippet

Yes. fadd of a zero is never something you want to introduce not on purpose, I suspect.

Looks like ReductionOp already has an optional accumulator. So its the lowering from vector.contract into this that needs to be fixed. Should be easy....

Yes. fadd of a zero is never something you want to introduce not on purpose, I suspect.

From HPC days my rule of thumb, dont rely on commutativity or associativity of floating-point adds. Just make sure you are doing identical left to right addition when you can (essentially treat it as a scan)

https://github.com/MaheshRavishankar/llvm-project/commit/fe926195cf7119fd7bc41d53a6eca9f060db4a93 fixes the issue in the dispatch that Hanhan had above. The output with and without host features flag is the same now
$ ~/iree/build_debug/tools/iree-run-module --flagfile=flagfile --module_file=repro.vmfb
EXEC @forward_dispatch_24
result[0]: hal.buffer_view
16xf32=2056.356689 996.230591 1129.984497 847.100647 1054.056274 1109.781860 834.492554 891.951904 2034.976929 81371.257812 1890.394165 1775.567017 2241.177734 2055.167480 2253.343506 2316.947021

It is different from both of what Hanhan saw earlier but that makes sense because even without host flag, things are still vectorized.

Nice, 81371.257812 is also the output of scalar implementation. This is the result we're looking for!

Thanks for persevering, folks (I know not landed yet, but root cause/verification is 90% of the battle)!

I am not sure we have all the root cause. Trying the models.

With Mahesh's patch, the e2e results of albert_base_v2 are identical between w/ and w/o host features. The max diff is within 1e-6. It indeed fixes an issue in IREE. However, there still are numerical issues between IREE and golden values (i.e., the values in expected_out.txt). One example is at 31-th number -- 0.058058 vs 0.0490268 where the max diff is 0.0178. I think there are other issues, I'll try to find them out.

@hanhanW I am not too sure about the expected_output.txt. If they are generated from TF, they might have TF related reassociation errors as well. Since IREE w/ and w/o host features are the same, that is less chance of a vectorization issue.

Great work. Is this patch something we can review and land while we continue to debug?

sending out for review now

Good point.. @powderluv is there a way to get the golden result? Can you attach them to the issue?

It is the assert at the top of this issue ?

I think that's from tensorflow, and might has reassociation errors as well? Based on the experiments above, I found that the reassociation errors might not small. Can you help verify if they are the golden result we'd like to compare with?

I ran several setup (with Mahesh's patch). I could not find suspicious things in IREE.

The results of scalar implementation (ie. use CPUDefault for all the kernels) and default configuration are identical. They are exactly the same.
IREE w/ and w/o host features are the same witinh atol=1e-6.
The VMVX and CPU backend are the same within atol=1e-3, rtol=1e-4.
vulkan and CPU results are match within atol=1e-4, rtol=1e-4.

hmmm...looks like the expected_out.txt in gs://iree-shared-files/nod-perf/ean/albert_base_v2_temps/ is wrong. I generated the expected_out.txt myself and they match with rtol=1e-02, atol=1e-03 config. The albert-base model looks good to me.
I'll look into other models.

Most of the expected_out.txt are quite wrong. I regenerated the files and tested them. Here is the update with Mahesh's patch.
Passing models with rtol=1e-2 and atol=1e-3 config:
albert-base-v2_test.py::AlbertBaseModuleTest::test_module_static_cpu
-- it's already passing

camembert-base_test.py::CamemBertModuleTest::test_module_static_cpu
max_diff = 6.53e-05

distilbert-base-uncased_test.py::DistilBertModuleTest::test_module_static_cpu
max_diff = 6.827e-5

mobilebert-uncased_test.py::MobileBertModuleTest::test_module_static_cpu
max_diff = 4.34e-5

mpnet-base_test.py::MpNetModuleTest::test_module_static_cpu
max_diff = 9.18e-5

roberta-base_test.py::RobertaModuleTest::test_module_static_cpu
max_diff = 6.301e-5

Failing models:
layoutlm-base-uncased_test.py::LayoutLmModuleTest::test_module_static_cpu
Failed @ position 0 : -1.532623 v.s. 1.50368261

The results of layoutlm are very different. I'll look into the model.

Most of the expected_out.txt are quite wrong.

That's weird, those really shouldn't be wrong.
Is this with the a2_highgpu instance?

That's weird, those really shouldn't be wrong.
Is this with the a2_highgpu instance?
@powderluv to check the fixes here.

They are not generated by IREE, so it doesn't matter which instance is used. I think the golden values should be identical. The fixes are in, please help check if it's still an issue on your side. I'll close the issue since I verified that it's fixed.

@hanhanW should we open a separate issue for the model that is still failing?

Yes, I'll open a separate issue for the model.

They are not generated by IREE, so it doesn't matter which instance is used. I think the golden values should be identical.
@hanhanW

The golden values from TF on the a2-highgpu-1g instance are different from the ones generated on other configs. Note that we've experienced this numerics issue only on the a2-highgpu-1g instances.
It's also worth noting that TF appears to utilize the A100 gpu to generate the model and golden values for our tests even when we are using CPU to generate mlir/vmfb and execute via IREE.
We're still experiencing the numerics issue for our a2-highgpu instances on latest IREE. Let me know if a fresh set of artifacts would be helpful.

The golden values from TF on the a2-highgpu-1g instance are different from the ones generated on other configs.

It sounds really bad. The golden values should be identical. They should not be platform/target dependent. It could let us run into wrong direction in terms of optimization. Because we won't really know if our optimization is acceptable to the model or not.
I can give it a shot on a2-highgpu-1g again if you can send us a fresh set of artifacts. Thanks!

Is this issue really closed? I had trouble following the last part about golden values. Seemed not our problem?

The issue about reassoc is fixed. The output values from IREE are accurate on my workstation. I haven't tested it on a2-highgpu-1g yet.
I'll check if IREE's results between different machines match or not later today. If they match, the issues are about "golden values". In this context, it's not IREE's problem.

I verified that the results from a2-highgpu-1g and my workstation are identical. The max diff is exactly 0. This is not a numeric issue on IREE side. The issue happens in "golden values". Please make sure that the TF implementation generates good outputs or compare the results with real golden values.

I filed a separate bug for tracking one of the associative reordering issues. I think this issue can be closed. One of the models still seems to have some correctness issue, that also probably needs to be tracked by a separate bug.
@powderluv I'd consider this particular issue closed. Can we run the crank again and file separate issue for any new issues?

One of the models still seems to have some correctness issue, that also probably needs to be tracked by a separate bug.

It's been tracked in https://github.com/iree-org/iree/issues/9716

authors: MaheshRavishankar ThomasRaoux benvanik hanhanW monorimet powderluv stellaraccident

repo: google/iree


text:
Update LLVM & TF Submodules.

Updates TF to HEAD and LLVM to match TF commit.
./scripts/git/update_tf_llvm_submodules.py --update_build_files=true
This is likely going to break the GitHub action bazel builds because of https://github.com/google/iree/issues/1737, but currently all the builds are broken because of the LLVM update, so we don't have much choice. Also the actual TF integrations are broken by TF and LLVM being out of sync at the moment.

Oh looks like the new llvm hash hasn't made its way to TF yet, actually

Updated again now that TF has exported the new hash

Failure is expected test failure from unidirectional_lstm that I'm trying to fix in https://github.com/google/iree/issues/1754

source: gharchive/pull-request

license_type: permissive

license_source: github-api


text:
Adding -verify-each flag to iree-run-mlir matching iree-opt.

Disabling the verifier with -verify-each=0 dramatically speeds up large
models.

It sounds obvious, but checking my understanding:
Verifiers have no side effects, so disabling them is like disabling warnings or logging? For tests and debug work, keep them on, but for production, disable them if you want extra performance?

yep! that's the idea - since iree-run-mlir is a tool meant for IREE developers (or people who are debugging IREE issues) it seems like a good idea to keep it on. Except for bert - that thing is too big :)

You made me think and I added it to iree-translate as well, so now iree-opt, iree-translate, and iree-run-mlir all accept it.


text:
Adding experimental synchronous executor using inline command buffers.

This is mostly just the plumbing to enable #4680 by defining a new dylib-sync runtime driver (pending #4298 to cleanup drivers).
This driver does not support recording/replaying command buffers and has some gotchas with semaphores (no deadline waits, probably corner cases for things not covered in the CTS, etc). It should get us something that builds easier on tiny systems, though.
This uses iree/base/synchronization.h for the semaphore implementation. The intent is that for bare-metal systems we can implement a version of iree/base/synchronization.h observing the specific requirements of the platform while still preserving the semaphore semantics - it's still useful to use this driver with systems with multiple threads (lots of dual-core MCUs/MPUs) or to enable pipelining. Systems that have no threads can no-op the mutex/notification (or disable interrupts/etc).
Trace of bert with -driver=dylib-sync looks pretty good:

(the large __init block at startup and free at shutdown will be solved by #5504 and #5472)
Note that due to the design here there are no changes needed in the compiler/runtime besides the single flag added in #5508 so the same exact module runs async as well:

That's pretty cool :)
Progress on #4680.

(typo'ed my own name, but github so it stays :)


text:
Closes: #1

updated validations.py python script.
Fixed the behavior of validate_user function in validations.py .
Fixes #<issue_number_goes_here>

It's a good idea to open an issue first for discussion.

[x] Tests pass
[x] Appropriate changes to README are included in PR

Fixed the behavior of validate_user function in validations.py
Closes: #1
Fixes <issue_number_goes_here>

It's a good idea to open an issue first for discussion.
-[] Test pass
-[] Appropriate changes to README are included in PR.

repo: google/it-cert-automation-practice


text:
Issue with SSLSocket

I've found a problem with j2objc's SSLSocket implementation. This time, luckily, it should be easy to fix:
In the current implementation, it's possible that the dealloc method throws an exception. This is due to the call to tearDownContext (which checks the ssl status and can throw an exception) here.
I guess the best solution would be to add a boolean parameter to tearDownContext that determines whether the checkStatus method should be called.
Please let me know what you think!

Good point! I think it should be enough to remove the call to checkStatus in tearDownContext. A "best-effort" approach should be OK given that we are already disposing the context.

+1
On Fri, Dec 7, 2018 at 8:31 AM Antonio notifications@github.com wrote:

Good point! I think it should be enough to remove the call to checkStatus
in tearDownContext. A "best-effort" approach should be OK given that we are
already disposing the context.
—
You are receiving this because you are subscribed to this thread.
Reply to this email directly, view it on GitHub
https://github.com/google/j2objc/issues/1016#issuecomment-445287800, or mute
the thread
https://github.com/notifications/unsubscribe-auth/AC2kr8dh4eTQlUqsK5VrvPvN-9-wOpGHks5u2pf7gaJpZM4ZIt3Z
.

Fixed in latest source.

source: gharchive/issue

authors: FD- antonio-cortes-perez tomball

license_source: bigquery


text:
The Package Prefixes page should mention that wildcard matching is possible

The Package Prefixes page here...
http://j2objc.org/docs/Package-Prefixes.html
... should mention that wildcard matching is possible with the --prefix and prefixes flags, i.e. com.company.*=AB.

Fixed, site should update in a few minutes.


text:
Question about block_until_ready() on tuple

I want to time the following:
opt_state = update(itr, grad(loss)(get_params(opt_state)), opt_state).
opt_state is a Python tuple so I can't call block_until_ready() directly.
What is the best way to ensure that opt_state is consumed from the host so I get accurate time?

nothing; does containment in a native Python contain imply the values have already been consumed?
tree_map and call block_until_ready() over all the leaves of opt_state
make opt_state a JAX type and call block_until_ready() once (If so, how to convert it to JAX type?)
directly consume from the host in some other way?

I think tree-mapping block_until_ready is a decent idea. I don't think it should add noticeable overheads (based on my guess about how much time the computation itself will take).

nothing; does containment in a native Python contain imply the values have already been consumed?

No, loops won't do anything special. The only thing that blocks the Python thread (e.g. so that timers are accurate) is executing a non-jax operation on it (like printing a value, which will entail blocking until that value is ready and then also transferring it to the CPU) or block_until_ready.

make opt_state a JAX type and call block_until_ready() once (If so, how to convert it to JAX type?)

We used to have JaxTuples! But they make the system much more complex, both in terms of "front-end" transformation stuff and "back-end" low-level runtime stuff.

directly consume from the host in some other way?

That works, e.g. printing the values, but then you'd also be timing the transfer-to-host time as well as whatever operation (e.g. printing) is being performed.
So yeah I'm thinking tree_map(lambda x: x.block_until_ready, opt_state)! But also if update is jitted then you can just do tree_flatten(opt_state)[0][0].block_until_ready(), since all results of a jitted function become available at the same time.

Thanks for the very detailed reply as always @mattjj :)

No, loops won't do anything special. The only thing that blocks the Python thread (e.g. so that timers are accurate) is executing a non-jax operation on it

Interesting, good to know!

We used to have JaxTuples! But they make the system much more complex, both in terms of "front-end" transformation stuff and "back-end" low-level runtime stuff.

Haha so I'm not crazy, I remember noticing these before I think! The way JAX handles nested containers is super nice. I suppose it's one of the simpler features but honestly one of my favourite things about JAX btw.

That works, e.g. printing the values, but then you'd also be timing the transfer-to-host time as well as whatever operation (e.g. printing) is being performed.

Good point, I guess that's why block_until_ready() is useful in the first place.

So yeah I'm thinking tree_map(lambda x: x.block_until_ready, opt_state)! But also if update is jitted then you can just do tree_flatten(opt_state)[0][0].block_until_ready(), since all results of a jitted function become available at the same time.

Ah, yes update is jitted so I think this is what I'll go with, thanks for pointing out this additional simplification.

authors: jacobjinkelly mattjj

repo: google/jax


text:
Deprecate X64 flag

This is the external tracking bug for deprecation of the X64 flag.
The final desired outcome here is the following:

X64 flag no longer exists
jax.numpy functions (e.g. jnp.arange, jnp.zeros, etc.) default to 32-bit, unless an explicit dtype is requested.
Python scalars default to weak-typed 64-bit

The plan is roughly the following:

Change jax.numpy function defaults to 32-bit. This should only affect users of X64 mode, and will require adding explicit 64-bit dtypes to some code. At this point, enabling X64 will produce the desired long-term behavior.
Change the default X64 flag state to True. Users can still recover X32 behavior by explicitly setting it to False
Deprecate setting the X64 flag
Remove the X64 flag

The main rationale is that the modalism in the API complicates JAX both from a user perspective and a developer perspective.
For the user, there's no way to reason about how a particular block of code will behave without global knowledge of the environment in which it is run (e.g. the value of the JAX_ENABLE_X64 environment variable at the time that JAX was imported).
Additionally, due to the way JAX tracing works, attempts to allow local overrides of the global state face fundamental challenges (e.g. https://github.com/google/jax/issues/5982).
Confusion about the x64 flag is one of the most common drivers of user confusion in our feedback forums.
Other similar frameworks (e.g. pytorch and tensorflow) do not have any such global modalism.
On the developer side, the global flag doubles the amount of CI testing we need to cover the package, and means that every line of code in the package must be written with this modalism in mind.
These are some of the reasons we've decided to explore the possibility of deprecating the X64 flag. I'm curious to hear your thoughts on those reasons! Thanks

May I ask in terms of specific examples?
How would jax.numpy.linalg.svd behave?  Would it always do the computation in 64-bit or would it depend on whether the type of the argument a is 64-bit float or 32-bit float? Or would the implementation internally increase precision to give accurate results even though it finally converts the output singular values to the data-type of a?
Several functions use an internal helper function _promote_arg_dtypes  to promote to an inexact type if the input is an integer type.  Would these functions then promote to 32-bit or 64-bit?
CPUs usually have 64-bit floating-point units by default. So 32-bit floating-point computation doesn't lead to computational savings? NumPy usually goes for 64-bit computation by default. Would JAX become more aligned with it or would it have different precision management rules?

The general idea is that JAX would do computations in the precision of the input type. In the presence of mixed inputs, JAX would decide the output type based on its existing Promotion rules.

Hi @jakevdp, I can see that this is a work in progress, so potentially annoying of me to ask this, but is it currently possible to emulate the behavior of
JAX_ENABLE_X64=True JAX_DEFAULT_DTYPE_BITS=32

without using environment variables? I am working on a small extension to a larger code-base, where I only need 64-bit dtypes in the extension and want to (as far as possible) preserve the behaviour of the rest, which was written with x64 disabled.
When I run
from jax.config import config
config.update("jax_enable_x64", True)
config.update("jax_default_dtype_bits", 32)

assert jnp.array([1.], dtype='float64').dtype == jnp.float64
assert jnp.array([1.]).dtype                  == jnp.float32

I currently get an assertion error from the final line, i.e. the config.update("jax_default_dtype_bits", 32) doesn't seem to have had any effect. Running the same snippet with the above environment variables set does work, but I'd like to avoid having to use environment variables if possible.
I also don't want to rely on API that is about to be removed, so if there is a better overall approach to this then please let me know.

Unfortunately the effect of the flag happens at import time, so changing the value after import will have no effect.
If you want to avoid setting system variables outside your Python process, you could do something like this:
import os
os.system['JAX_DEFAULT_DTYPE_BITS'] = '32'

import jax

So long as jax is not imported until after the system variable is set, it should work.

Thanks, that should do the trick for now.

@jakevdp do you still plan to deprecate the flag, or you changed plan and this likely won't happen anymore?

This has proven difficult to deprecate, because so many users rely on the current behavior. I'm going to leave this open, but de-prioritize it.

I think we've come to the conclusion that we can't remove the X64 flag, given the extent that users depend on it. I'm going to close this.

authors: PhilipVinc j-towns jakevdp shailesh1729


text:
[jax2tf] Added a flag and environment variable to control the serialization version

This allows us to control the serialization version to be compatible with
the deployed version of tf.XlaCallModule. In particular, we can run
most tests with the maximum available version, while keeping the
default lower.

@junwhanahn @rxsang PTAL

The changes look good to me.
Somewhat related question: do you plan to maintain min/max supported version in jax2tf (if we ever need one when we clean up some old jax2tf logic) separately from TF2XLA, or in sync? I would guess that the former makes more sense, but not sure if it will make the testing more complex.

I don't quite see how to keep the versioning perfectly in sync. My plan is to support in jax2tf the versions that have been active in google3 over the last month. This means 1 or 2 versions, the most recent. The tests in google3 will verify that we can use the most recent version and the one that is marked as default. There is no test per se that we can use the version that has been active a month ago; we rely on people not updating the default serialization version too early.
Also in OSS people will have different versions of JAX and TF. I do assume that people will be willing to use tf-nightly if they get an error; this has been true for jax2tf so far because it relies always on some very recent details for the TF ops.


text:
Add a cargo command for kythe [phab:D1010]

Created by djrenren at 2016-08-13 00:48:57:
This adds an indexer driver for cargo so that one could write:

cargo kythe for a top-level index

cargo kythe full-index for a deep index including dependencies

cargo kythe web to run the web UI

The utility requires a KYTHE_DIR environment variable to set and complains if it is not present.

zarko wrote at 2016-08-15 16:55:09:
Please add a README with instructions on how to build and run the tool.
Bazel complains when I tried to do the obvious thing:
bazel build //kythe/rust/cargo-kythe
..
WARNING: kythe-trunk/kythe/rust/cargo-kythe/BUILD:5:1: target 'cargo-kythe' is both a rule and a file; please choose another name for the rule.
INFO: Found 1 target...
ERROR: kythe-trunk/kythe/rust/cargo-kythe/BUILD:5:1: Executing genrule //kythe/rust/cargo-kythe:cargo-kythe failed: bash failed: error executing command /bin/bash -c ... (remaining 1 argument(s) skipped): com.google.devtools.build.lib.shell.BadExitStatusException: Process exited with status 101.
    Updating registry `https://github.com/rust-lang/crates.io-index`
   Compiling rustc-serialize v0.3.19
   Compiling indexer v0.1.0 (file://kythe-trunk/kythe/rust/indexer)
   Compiling cargo-kythe v0.1.0 (file://kythe-trunk/kythe/rust/cargo-kythe)
error: failed to run custom build command for `cargo-kythe v0.1.0 (file://kythe-trunk/kythe/rust/cargo-kythe)`
process didn't exit successfully: `kythe-trunk/kythe/rust/cargo-kythe/target/debug/build/cargo-kythe-b4d6f68a7c42e6f5/build-script-build` (exit code: 101)
--- stderr
thread 'main' panicked at 'called `Result::unwrap()` on an `Err` value: NotPresent', ../src/libcore/result.rs:788
note: Run with `RUST_BACKTRACE=1` for a backtrace.

Target //kythe/rust/cargo-kythe:cargo-kythe failed to build
Use --verbose_failures to see the command lines of failed build steps.
INFO: Elapsed time: 25.990s, Critical Path: 20.42s

zarko commented inline (context lost during migration) at 2016-08-15 16:56:33:
nit: whitespace

djrenren wrote at 2016-08-15 17:43:43:

Add expect messages to build.rs

djrenren wrote at 2016-08-15 22:54:04:

Prep for publishing. Remove from bazel build path

zarko commented inline (context lost during migration) at 2016-08-15 23:02:41:
nit: cargo-kythe

zarko commented inline (context lost during migration) at 2016-08-15 23:02:41:
nit: use cargo (e.g., the monospaced command name) or Cargo (the proper name of the tool).

zarko commented inline (context lost during migration) at 2016-08-15 23:02:41:
Where are the index data written and how can that be configured?

zarko commented inline (context lost during migration) at 2016-08-15 23:02:41:
nit: you probably want "dependencies" here, since "dependents" goes the other way.

djrenren wrote at 2016-08-16 17:59:47:

Fix README nits

zarko wrote at 2016-08-16 19:47:39:
Before you publish this to crates.io (or elsewhere), we need to package a new Kythe binary release that includes #1339.

zarko accepted this revision at 2016-08-16 19:47:39

djrenren closed via commit 0a4e1e8021309abb8f377dd09bf12687711b2f10 at 2016-08-16 21:00:09

repo: google/kythe


text:
Add graphs integration tests for pythonboard

Adds integration tests for accumulating-context and math to pythonboard

It's a good idea to open an issue first for discussion.

[ ] Tests pass
[ ] Appropriate changes to documentation are included in the PR

Some work I had done a few weeks before my trip.
This runs integration tests for graphs in python. This change introduces simple ones like accumulating-context and math, but I'd like to add more complex graphs into the suite in follow-up commits.

repo: google/labs-prototypes


text:
Error if directory we're trying to create the database in doesn't exist

Original issue 167 created by kkooporation on 2013-05-11T09:57:47.000Z:
What steps will reproduce the problem?

open a database
bazinga!

(we should ditch the standard questions here..)
/What is the expected output? What do you see instead?
What version of the product are you using? On what operating system?
Please provide any additional information below./
When a database is being opened in a folder which itself does not exist, LevelDB errors with a dubious error message:
OpenError: IO error: ./db/root.db/LOCK: No such file or directory
at [..]/levelup.js:113:25
As this came up quite some times, we discussed the process here; either the path will recursively be created or instead an error being thrown indicating the specific issue.
We've been talking about this issue on GitHub regarding LevelDOWN / UP [the NodeJS bindings for LDB] (https://github.com/rvagg/node-levelup/issues/135) and came to the conclusion this would have to be implemented in LevelDB itself.
Of course, I'd welcome a discussion here — thanks!

you shall use absolute path, because the shell+leveldb may can't figure out  "./" and "~/" relative path correctly.

:+1:

Agree with commentors - using relative paths (which are platform specific) is not a good idea.

authors: cmumford hgqislub ralphtheninja

repo: google/leveldb


text:
Ninja : Error

When i am running the project in android studio i got the error as :
Build command failed.
Error while executing process E:\Android\Studio\cmake\3.6.4111459\bin\cmake.exe with arguments {--build C:\Users\rajarajanayyarasu\Desktop\live-transcribe-speech-engine-master\live-transcribe-speech-engine-master\app.externalNativeBuild\cmake\arm7Debug\armeabi-v7a --target ogg_opus_encoder_tool}
ninja: error: '../../../../third_party/libopus/lib/armeabi-v7a/libopus.a', needed by '../../../../build/intermediates/cmake/arm7/debug/obj/armeabi-v7a/libogg_opus_encoder_tool.so', missing and no known rule to make it
Please Give some solutions to fix it.I dont know how to fix it.

hi, could you provide the "build_log.txt" for checking. Thanks.

Hi,
The failed is related with build libraries out using our sample Android application.
Currently, these instructions assume that the host operating system is an Ubuntu-like flavor of Linux. For other operating systems likes as Window, we cannot promise that run well in other build environments.
However, it is an optional solution, you can refer attached patch file which striped out the native libs of compressing audio.
You don't need to run build_all.sh to generate APK, instead you could build it by your android studio.
0001-Remove-Opus-native-libs-for-windows.patch.zip

repo: google/live-transcribe-speech-engine


text:
Fix missing SSE detection on x64 targets. Fixes #25

One important note on this PR: with this change, matrix_benchmarks takes twice as long to run on my test system with SIMD enabled as it does with SIMD disabled. That seems... unintuitive. And worth investigating further.

Re the CLA: I'm a Google employee. I just registered my GitHub account, though, so approval may still be processing.

To be precise: It takes twice as long to run for x64, and we don't know how much slower it was before, since on x64 it wasn't using SIMD at all, right?
This change seems otherwise ok to me, it is not like anything in this PR is causing a slowdown by itself, so I think it should go in. Finding out the slowdown is a separate issue.

authors: cdwfs gwvo

repo: google/mathfu


text:
how can I enable/disable the red rectangle of face detection in android?

I want to take a picture when detects the face, now I use screenshot to implement.
But the result contain the "face, 0.99" and some point and a rectangle.
It's possible to close it?
The camera will show like below.

I need to get a picture like below.

@gkarpiak @mcclanahoochie
Can anyone help me? thank you.

the rendering happens in the AnnotationOverlayCalculator , you can replace this node with a PassThroughCalculator like this:
node {
  calculator: "PassThroughCalculator"
  input_stream: "throttled_input_video"
  output_stream: "output_video"
}

or, you can edit the DetectionsToRenderDataCalculator to output empty RenderData packets.

Thanks for your help!! @mcclanahoochie @sgowroji
Finally I resolve the problem, let me talk about my solution in detail!

Edit the following file and rebuild face detection aar.
/mediapipe/mediapipe/graphs/face_detection/BUILD
cc_library(         name = "mobile_calculators",         deps = [             "//mediapipe/calculators/core:pass_through_calculator"    <- Add this line         ],     )
/mediapipe/mediapipe/graphs/face_detection/face_detection_mobile_gpu.pbtxt
node {       calculator: "PassThroughCalculator"       input_stream: "throttled_input_video"       output_stream: "output_video"   }

After rebuilding mobile_gpu.binarypb and changing the file mentioned above.

authors: mcclanahoochie solmate-james

repo: google/mediapipe


text:
Android build crashes when model_complexity set to 0 in posetrackinggpu

Issue:
When building with model_complexity set to 0 (through variable select: 0) in pose_landmark_model_loader.pbtxt  the app crashes. It works fine when the same is default of 1. I'm simply running the sample app with only this change.
Platform: Android Mobile
Errors: No error while building, app apk starts and crashes.

Hi @Raghav997,
This can be an interesting finding. Would you let us know the complete steps followed to build the aar to understand the issue better. We suggest to fill the template to get maximum information to debug on the issue.

I just followed basic instructions given here.
[(https://google.github.io/mediapipe/getting_started/android.html)]
To replicate, just follow the instructions. You'll see the "1" model is working whereeas the lite "0" is crashing.

Any update on this?

Hi @Raghav997,
It looks like the two of one Tflite model files being missing from build file
"third_party/mediapipe/modules/pose_landmark/pose_landmark_lite.tflite" missing.
"third_party/mediapipe/modules/pose_landmark/pose_landmark_full.tflite" found.
Could you verify that "pose_landmark_lite" is actually missing from the BUILD file? If missing, you might need to add the dependency to pose_landmark_lite.tflite to the BUILD file.

In Ubuntu, I just downloaded the mediapipe git source ($ git clone https://github.com/google/mediapipe.git) and build the holistictrackinggpu android solution.
$ bazel build -c opt --config=android_arm64 mediapipe/examples/android/src/java/com/google/mediapipe/apps/holistictrackinggpu:holistictrackinggpu
and the holistictrackinggpu.apk (which was built in the bazel-bin/mediapipe/examples/android/src/java/com/google/mediapipe/apps/holistictrackinggpu/holistictrackinggpu.apk) is working well on my Android Phone.
Do you know how to set 0 for the value "model_complexity" for holistictrackinggpu Android ?
I just modified the mediapipe/mediapipe/modules/pose_landmark/pose_landmark_model_loader.pbtxt  as like below.
node {
calculator: "SwitchContainer"
input_side_packet: "SELECT:model_complexity"
output_side_packet: "PACKET:model_path"
options: {
[mediapipe.SwitchContainerOptions.ext] {
select: 1
and built again ($ bazel build -c opt --config=android_arm64 mediapipe/examples/android/src/java/com/google/mediapipe/apps/holistictrackinggpu:holistictrackinggpu)
but the holistictrackinggpu.apk is not working on the Android Phone.
I want to set 0 for the value "model_complexity" for holistictrackinggpu android.

authors: Raghav997 kuaashish taeyeonlee


text:
xr-environment attribute causes errors the second time a webxr session is started

Description
When using xr-environment attribute, reopening a webxr session (after having closed a previous one)
causes NPE errors in the console
Live Demo

https://modelviewer.dev/examples/augmentedreality/#ar

Version

model-viewer: v3.4.0

Browser Affected

[x] Chrome, version: 117
[ ] Edge
[ ] Firefox
[ ] IE
[ ] Safari

OS

[x] Android
[ ] iOS
[ ] Linux
[ ] MacOS
[ ] Windows

AR

[x] WebXR
[ ] SceneViewer
[ ] QuickLook

@bialpio @toji Can you repro this on your devices?

I can't repro on my devices. Also, I'll note that Chrome 117 is fairly old, having been released 7 months ago. The current release is 124. If @LR17 can update to the latest version and try again, I'd appreciate knowing if it still repros for them!

Yes, same error, also with chrome 123

Thanks for confirming! I tried it again this morning and I still can't provoke the issue that you're talking about. Could you also provide the device that you're testing with and also the version of ARCore that's installed on your device? (You can find this by going to Settings->Apps->See All->Google Play Services for AR. The version is listed at the bottom.)
Also, when this error appears does it cause a noticeable disruption to the AR content? (Black screen, no model shown, corrupted display, etc?) Or does the model display as expected and the only indication of a problem is the console messages?

Thank you again! I really appreciate your willingness to follow up on this.
I'm guessing this is the device that you're referring to. If so, then it should be an ARCore supported device, according to the official support list. There's no compatibility notes listed there indicating that lighting estimation may not work or anything like that. The GPU is relatively low-end so it's possible that it's just not able to keep up with lighting estimation and AR tracking at the same time, but I'd sort of expect there to be a note to that effect it that were the case.
As for the ARCore version, I have a slightly newer one on my devices (1.42) so you may want to see if there's any updates available for "Google Play Services for AR" in the Play store, but it's not too old so I'd be somewhat surprised if that was contributing.
@elalish: I'm going to guess that this particular device is getting into a state where it's either not providing valid lighting estimates or providing them at a lower/more erratic rate, and that's throwing off either Chrome's implementation or Three's/model-viewer's expectations. Worth looking at the Three code first to see if there's any bad assumptions being made about consistency of lighting estimation data being provided. (If so it's probably my fault, I think Three is still using some of my code for that.)

authors: LR17 elalish toji

repo: google/model-viewer


text:
Add a scrap command to backfill Spec11 threats

This change is

ptal?

repo: google/nomulus


text:
Ignore instrumentation events from threads spawned by injected lib.

Directly after the injection of the library we create a new thread using the
clone syscall and execute the initialization code in this thread.
We resume the execution of the game and allow the initlization to happen. After
that we detect the newly spawned threads, attach again to the target and set
the tids in the lib so that we can ignore these threads in EntryPayload.
The bug contains some additional context.
Bug: b/236121842
Test: Covered by existing unit tests. Also manual testing with several
targets.

Oooh right. So the actual question was: what if the clone syscall fails (returning -1 in rax for the "parent"). I guess the timeout in GetNewOrbitThreads takes care of it?

Yes, this is a bit lazy. I didn't want to implement the error handling in assembly / machine code. The error we get from GetNewOrbitThreads is close enough I think.

repo: google/orbit


text:
JVM: Fix wrong query parameter bug

The introspector query for jvm return type searching has a wrong parameter name. This PR fixes that.

/gcbrun skip

repo: google/oss-fuzz-gen


text:
[UI] Mousing over counter tracks set to "Rate" displays incorrect values

If i mouse over counter tracks that are set to "rate" in the drop down, it seems to print the raw values rather than the rate values.
For the power_rails.pb example below this means the mouse over is showing energy instead of power.
How to reproduce

open power_rails.pb example:

./tools/open_trace_in_ui -i test/data/power_rails.pb

mouse over the values in one of the counter tracks.

Observed behavior

as we move the cursor along the timeline you can see the value continuously increate
if toggle the track to "zero-based" we notice the mouse over has the same value"

https://github.com/google/perfetto/assets/77857237/bb77bb4d-2335-4de6-975a-5a26932d16c6
Expected behavior
because the mouse over dialog says "delta/t" it is expected that we're printing a change in energy over time (AKA power)
we expect the value printed in the mouse over to not to continuously increase and instead go up and down according to the magnitude of the chart displayed on the timeline
The values should be similar to the "Rate /s" that is displayed if you click and drag over a section of the trace:

Took a quick look since this was sitting idle for a few weeks.
looks to me like the code was missing a this.config.scale === 'RATE' check in one spot causing the symptoms above.
I think this should fix it: https://android-review.googlesource.com/c/platform/external/perfetto/+/2744144

Great! Thanks a lot for fixing that 👍
It's been merged and should be on the tip of main/autopush shortly.

repo: google/perfetto


text:
Could a zip of the UI static site build be added as an artifact to releases?

The site's UI is built using NPM etc, which can be done in a container as described here (https://github.com/google/perfetto/issues/50).
But it would be very convenient if the UI build artifact could be uploaded as a file in the release files too.
Could this be done for future releases please?
Thanks for all your work on such a comprehensive tool.

It's quite easy to just build the UI and it would be a bunch of work for us to integrate into all our automatic release tooling so I doubt we'll get to this any time soon.

authors: LalitMaganti inbilla


text:
Halt() does not unblock WaitForEdge()

According to the documentation, calling Halt() or In() on a pin should cause WaitForEdge() to unblock and return false. When I attempt to do this on a Raspberry Pi Zero W, however, WaitForEdge() does not become unblocked.
package main

import (
	"fmt"
	"time"

	"periph.io/x/periph/conn/gpio"
	"periph.io/x/periph/conn/gpio/gpioreg"
	"periph.io/x/periph/host"
)

func main() {
	if _, err := host.Init(); err != nil {
		panic(err)
	}

	aPin := gpioreg.ByName("23")
	aPin.In(gpio.PullUp, gpio.BothEdges)

	go func() {
		time.Sleep(1 * time.Second)

		fmt.Println("calling Halt()")
		aPin.Halt()
		fmt.Println("called Halt()")
	}()

	fmt.Println("wait for edge")
	aPin.WaitForEdge(-1)
	fmt.Println("done waiting for edge")
}

This program will output:
wait for edge
calling Halt()
called Halt()

and then block indefinitely.
Platform:

OS: Raspbian GNU/Linux 9
Board: Raspberry Pi Zero Wireless

Sorry I didn't have time to play with this locally and won't likely have time until this weekend. :/

Sorry I didn't have time to play with this locally and won't likely have time until this weekend. :/

Don't worry about it. I don't need it for anything critical -- it's only for a small side project I'm working on.

Manually ported to https://github.com/periph/host/issues/36.

authors: maruel toxygene

repo: google/periph


text:
Address bar

I propose adding an address bar to the preview window, allowing the user to inspect and change the iframe's url. For example if the playground app is a multipage or singlepage app with a router

I think it's a good idea!

Another great reason to have opt in support for an address bar is for demo-ing labs/router ;)

authors: AndrewJakubowicz aomarks bennypowers

repo: google/playground-elements


text:
Improve the document in detail

Dear teachers
I am  a follower of this project.  I have a problem recently.  I read the chapter of Key Derivation , but I don't quite understand how to generate secret keys for a psp stream. Could you please tell me the detail of initial handshaking and key exchange？

Use DTLS ？

The shared key generation via the KDF is documented and straight forward. I think what you might be asking is how to get that shared secret key generated on the receiver back to the sender. Likely, this is a job for TLS borrowing its handshake. This is what QUIC does and implementations hook into the handshake function callback mechanisms found in TLS libraries.

If we ignore the details of key exchange and just assume an out-of-band channel for communication between the sender and receiver, the steps would be:

Sender requests a new derived key from receiver.
Receiver picks a new SPI value and uses it to generate a derived key from the master key.
Receiver sends SPI and derived key to sender, which the sender can use to produce PSP-encapsulated packets.

authors: hedong0411 hlrichardson insanum

repo: google/psp


text:
Typo Error

A typo error in the comment of the apply_directed_instr() function of the riscv_instr_test_lib.sv file.
virtual function void apply_directed_instr();   // Mix below directed instructino streams with the random instructions     asm_gen.add_directed_instr_stream("riscv_load_store_rand_instr_stream", 4);     asm_gen.add_directed_instr_stream("riscv_loop_instr", 3);     asm_gen.add_directed_instr_stream("riscv_jal_instr", 4);     asm_gen.add_directed_instr_stream("riscv_hazard_instr_stream", 4);     asm_gen.add_directed_instr_stream("riscv_load_store_hazard_instr_stream", 4);     asm_gen.add_directed_instr_stream("riscv_multi_page_load_store_instr_stream", 4);     asm_gen.add_directed_instr_stream("riscv_mem_region_stress_test", 4);   endfunction

Thanks for fixing the issue

repo: google/riscv-dv


text:
Should the builder cache types?

For example:
    let void = b.type_void();
    let void1 = b.type_void();

results in
%2 = OpTypeVoid
%3 = OpTypeVoid

I expected just:
%2 = OpTypeVoid

Should the builder cache types? Or should duplicate instructions be removed in a later stage?

The SPIR-V spec requires that "Non-structure types (scalars, vectors, arrays, etc.) with the same operand parameterization cannot be type aliases." So it is actually invalid to have two s for the void type.
However, it is arguable whether this kind of "semantic" validation checks should happen in mr::Builder. mr::Builder is designed to be low-level with little sanity check. It is intended to be used by some higher level builder actually. (Sorry about this maybe unnecessary layers and abstractions.) Besides, SPIR-V spec 2.16 contains a bunch of validation rules. If we enable this one, what about the others?
I am planning a higher level builder (#5) and started writing code for it, but cannot find time for it right now. We'll definitely need a type cache there. ;-P

Okay sorry I haven't completely read the spec yet, I should probably do that now.

authors: MaikKlein antiagainst

repo: google/rspirv


text:
Update Plugin Manager to remove unneeded plugin dependencies

When a plugin is uninstalled the plugin manager should also remove all dependencies that are not needed for any other plugin. I have written code to do this for the "Remove Plugin and dependencies" plugin.

The issue is we need to somehow mark what plugins were manually installed.  It's entirely possible, for instance, that a manually installed plugin (Phoenix Core for instance), is meeting a prerequisite for only a single downstream plugin.  In most cases, you would not always want removing BMT to also remove the Phoenix Core.  Even though BMT might have been the only plugin that was using Phoenix Core, we can't assume that removing BMT should also remove the core.  right?

Correct.  I wrote some code that uses a recursive algorithm to walk the
plugin tree and only remove those dependency plugins that are not
dependencies of other plugins.  Take a look here:
https://code.google.com/p/tmiranda/source/browse/#svn%2Ftrunk%2FRPAD%2Fsrc%2Ftmiranda%2Frpad
Once I wrap my head around the core I could take a shot at it but right now
I'm not up to it.
On Mon, Aug 17, 2015 at 9:11 PM, Christopher Piper <notifications@github.com

wrote:

The issue is we need to somehow mark what plugins were manually installed.
It's entirely possible, for instance, that a manually installed plugin
(Phoenix Core for instance), is meeting a prerequisite for only a single
downstream plugin. In most cases, you would not always want removing BMT to
also remove the Phoenix Core. Even though BMT might have been the only
plugin that was using Phoenix Core, we can't assume that removing BMT
should also remove the core. right?
—
Reply to this email directly or view it on GitHub
https://github.com/google/sagetv/issues/18#issuecomment-132015281.

I agree with Fuzzy's comment...I don't think this is a good idea to do unless we are also tracking which plugins were manually installed. Otherwise you could remove something the user manually requested to install.

Maybe make it an option?  I think there are very few users who will
manually install a plugin that happens to be a dependency of another
plugin. Why would a user do that in the first place?
On Aug 18, 2015 2:38 PM, "Jeffrey Kardatzke" notifications@github.com
wrote:

I agree with Fuzzy's comment...I don't think this is a good idea to do
unless we are also tracking which plugins were manually installed.
Otherwise you could remove something the user manually requested to install.
—
Reply to this email directly or view it on GitHub
https://github.com/google/sagetv/issues/18#issuecomment-132312697.

Fuzzy gave a reason why someone would do it. :)  And I don't see that much actual benefit from removing plugins that are no longer needed as they aren't hurting anything....but if you remove something the user actually wanted there, then that's a problem.
Making some other option that removes all the unneeded dependencies via first figuring them out, and then prompting the user about what's going to be removed, that would be fine with me. :) But would of course then require extra things in the STV to make that happen. Lol...I guess this could actually be a plugin. :D

Ok, you win :)
On Aug 18, 2015 3:19 PM, "Jeffrey Kardatzke" notifications@github.com
wrote:

Fuzzy gave a reason why someone would do it. :) And I don't see that much
actual benefit from removing plugins that are no longer needed as they
aren't hurting anything....but if you remove something the user actually
wanted there, then that's a problem.
Making some other option that removes all the unneeded dependencies via
first figuring them out, and then prompting the user about what's going to
be removed, that would be fine with me. :) But would of course then require
extra things in the STV to make that happen. Lol...I guess this could
actually be a plugin. :D
—
Reply to this email directly or view it on GitHub
https://github.com/google/sagetv/issues/18#issuecomment-132323126.

What about the ability to see what a plugin is dependent on and what dependents on it in its properties.  That at least allows for manually removing unneeded plugins.

Not a bad idea at all. :)

authors: CraziFuzzy Narflex kkuhler tmiranda1962

repo: google/sagetv


text:
Increase LocInstWidth to 8 and add a assertion

The current LocInstWidth of 6 can be easily exceeded when using the whole component library with multiple width instances.

Google » souper #486 SUCCESS
This pull request looks good
(what's this?)

LGTM

authors: buildhive regehr rsas

repo: google/souper


text:
resolve: retain globals across REPL chunks

Previously, in the REPL, the globals bound by one chunk would become
the "predeclared" bindings of the next chunk. However, this produced
the wrong result for [x = x + 1] because both occurrences of x would
resolve to the same global x, which was undefined when the right-hand
side expression was evaluated, leading to a dynamic error. It would
fail for [x += 1] for a similar reason.
This change extends the resolver to allow clients to provide it with a
non-empty set of module globals. Thus the globals of one chunk remain
the globals of the next chunk, and the predeclared set is always empty.
The new API functions is intended for use only by the REPL.
Fixes #233

@laurentlb @brandjon

repo: google/starlark-go


text:
camera calibration

Which tool did you use to perform 'camera calibration'?

We used various approaches -- the right tool probably depends on your use case. If you have two cameras and can physically calibrate them, you could try using the Bouguet camera calibration toolkit: http://www.vision.caltech.edu/bouguetj/calib_doc/
Other possibilities would include running a structure from motion pipeline like COLMAP on your images (http://colmap.github.io).

authors: Songminkee snaves

repo: google/stereo-magnification


text:
prog: allow recursion for optional pointers

When syzkaller generates arg that uses a few structs that reference each other via pointers, it can go into infinite recursion and crash.
Fix this by forcing pointer args to be null when the depth of recursion reaches 3 for some struct.

LGTM

authors: dvyukov xairy

repo: google/syzkaller


text:
Compiling Issues on AArch64

Hi All,
I've got a compiling issue on one AArch 64 machine. The error message is
./tcmalloc/internal/percpu_tcmalloc.h:483:4: error: undefined named operand 'rseq_sig'
It seems that the RESQ support was added on Issue #50, but I am stilling get the compiling issues.
Does someone have met similar issues and be able to solve this?
Great Thanks!

Can you provide some more information about your compiler and which commit you're building from?

ubuntu20.04,clang-12 error.
ld.lld: error: undefined symbol: __rseq_abi

referenced by percpu_rseq_aarch64.S:217 (../../third_party/tcmalloc/tcmalloc/internal/percpu_rseq_aarch64.S:217)
percpu_rseq_asm.o:(TcmallocSlab_Internal_PerCpuCmpxchg64) in archive obj/third_party/tcmalloc/tcmalloc/internal/libpercpu.a
referenced by percpu_rseq_aarch64.S:217 (../../third_party/tcmalloc/tcmalloc/internal/percpu_rseq_aarch64.S:217)
percpu_rseq_asm.o:(TcmallocSlab_Internal_PerCpuCmpxchg64) in archive obj/third_party/tcmalloc/tcmalloc/internal/libpercpu.a
referenced by percpu_rseq_aarch64.S:217 (../../third_party/tcmalloc/tcmalloc/internal/percpu_rseq_aarch64.S:217)
percpu_rseq_asm.o:(TcmallocSlab_Internal_PerCpuCmpxchg64) in archive obj/third_party/tcmalloc/tcmalloc/internal/libpercpu.a
referenced 9 more times
clang-12: error: linker command failed with exit code 1 (use -v to see invocation)

@ckennelly

authors: Fllower ckennelly lrh12580

repo: google/tcmalloc


text:
Update/replace syntax error UI

Remove syntax error history and replace it with something that occupies less space (e.g., consider a "spell-checker" UI with squiggly lines under the problematic code and an error message displayed when user clicks on the code snippet).

Fixed in #384. Thanks @eledavi!

authors: rabidbit seanlip

repo: google/tie


text:
On Android, when GeneralSecurityException "decryption failed" is thrown when calling AeadFactory::decrypt?

Hi,
Edit: we are using Tink 1.2.2.
We have been getting crash reports with the exception GeneralSecurityException: decryption failed with the stack trace pointing to AeadFactory::decrypt as the source of the exception, and we are lost on when that happens.
We are not sharing the keys or the cipher text, so, the same device and algorithm that encrypts it, is the same one that is decrypting it.
We are also using the AndroidKeysetManager.
The error seems to happen exclusively on Android 8 e 9 so far.
Any idea on what may be causing this exception being thrown? Thank you :)
Check our full implementation: https://gist.github.com/AllanHasegawa/85431b6f7c53074511527655712c9872

We have been getting crash reports with the exception GeneralSecurityException: decryption failed with the stack trace pointing to AeadFactory::decrypt as the source of the exception, and we are lost on when that happens.

It happens when the ciphertext is invalid or the key is incorrect.

https://gist.github.com/AllanHasegawa/85431b6f7c53074511527655712c9872
Interesting.

Are you using the same keysetFileName? Can you dump the Shared Preference to make sure it is there and is not changed on each run?
Also: https://gist.github.com/AllanHasegawa/85431b6f7c53074511527655712c9872#file-tink_keystore-kt-L38 looks incorrect? You simply concatenate the associate value with ciphertext, so values.size should be 2?

It happens when the ciphertext is invalid or the key is incorrect.

hmmm. I'm saving the cipher text on SharedPreferences. Maybe Android is corrupting this data somehow? will investigate further.

Are you using the same keysetFileName?

Yes, same keysetFilename.

Can you dump the Shared Preference to make sure it is there and is not changed on each run?

Unfortunately I was not able to replicate this issue to know more about it. The crashes are coming from Crashlytics. I tried this code in many different devices and it always works.
On the devices that the code do work the contents of the SharedPreferences seems to be OK. I tried even manually changing the contents of the shared preference and even so I could not replicate the crash. If I change the key cache, I just lose the content, but no crash.
I'm running a test where it encrypt/decrypts billions of random strings with no issues. The problem should be an issue somewhere else.

values.size should be 2?

The regex matcher will return 3 groups. The first one is the full match, with the others 2 being the associated data with the cipher text. This works fine 99.97% of the time (0.03% being the crashes).

It looks like we've got a Heisenbug.

The error seems to happen exclusively on Android 8 e 9 so far.

Do you have a list of device models and/or manufacturers on which it happened?

Do you have a list of device models and/or manufacturers on which it happened?

Yes.

LG G7 ThinQ
Mi A2 Lite
Redmi Note 5 Pro
Motorola One
Moto Z3 Play
Moto E (5)
Galaxy A7 (2018)
Galaxy S9+
Galaxy Note9

If I read the code correctly, this would happen if clock.now().toString() returns a string containing a colon. It sounds unlikely this is the problem (why would it only happen so infrequently?), but still: is this possible?

I see.
Is it possible for you to do a "safety decryption" right after you encrypted the data? I.e., after
val cipherBytes = aead.encrypt(
plainText.toBA(),
associatedData.toBA())
add something like
val unused = aead.decrypt(cipherBytes, associatedData.toBA());
This would make the correct team is working on debugging this issue.

@tholenst Will do. Will create an exception to know if we are crashing right after generating the cipher text, or when restoring it from SharedPreferences.

We rolled out an updated version of this "safety decryption" as @tholenst suggested and we got no exceptions thrown when deciphering just after encrypting our data. This makes me think that Tink is working as intended here.
We are, however, still getting crashes when deciphering persisted data. But for now, I'll close this issue since I don't think there's much Tink can do to fix in this case.

I had a similar issue on a project that i'm working on. After stepping through the debugger to see what methods and classes get called, I noticed that from the Android integration package (https://github.com/google/tink/tree/master/java/src/main/java/com/google/crypto/tink/integration/android), AndroidKeystoreAesGcm.java does not get invoked but AesGcmJce does instead.
What I attempted to do was to use the static getOrGenerateNewAeadKey(String keyUri) in AndroidKeystoreKmsClient after generating the KeysetHandler like so (roughly based on your gist):
private fun getAEAD(): Aead {         generateKeysetHandler()         return getOrGenerateNewAeadKey(masterUri.toString())     }
This now invokes AndroidKeystoreAesGcm. My question to the team is: How should we properly obtain the Aead class for Android?

My apologies, I stand corrected. I've managed to replicate the scenario and in our case it occurs because when the user logs out, we delete the KeyStore alias entry. Unless the app is closed fully at that point, the next time a user logs in, it appears that even if we call generateKeysetHandler(), the subsequent attempt to get the KeysetHandle would fail.
Though my point that AndroidKeystoreAesGcm does not get invoked still holds true. Let me know if my usage above is correct.

Hi @daewin

Hey @thaidn. I believe that key generation etc. works as intended, but the problem was that we were lazy loading the Aead instance (getOrGenerateNewKeysetHandle().getPrimitive(Aead.class)), so even after deleting the KeyStore alias, it was still using the now defunct Aead instance on the next user login - the proper way is to generate a new Aead key and use that.
For some context, it looks a little like this:
    val aead: Aead by lazy {         AndroidKeysetManager.Builder()             .withSharedPref(applicationContext, TINK_KEYSET_NAME, PREF_FILE_NAME)             .withKeyTemplate(AeadKeyTemplates.AES256_GCM)             .withMasterKeyUri(MASTER_KEY_URI)             .build()             .keysetHandle             .getPrimitive(Aead::class.java)     }

Though i'm curious as to why AndroidKeystoreAesGcm never gets used even if i'm on Android APIs >=23. We seem to be using AesGcmJce exclusively.

Sorry for the late response. AndroidKeystoreAesGcm is only used to unwrap the keyset. AesGcmJce handles the actual data encryption.
Please reopen if you need more help.

Though i'm curious as to why AndroidKeystoreAesGcm never gets used even if i'm on Android APIs >=23. We seem to be using AesGcmJce exclusively.

Sorry for the late response. AndroidKeystoreAesGcm is only used to unwrap the keyset. AesGcmJce handles the actual data encryption.
Please reopen if you need more help.

So this makes some more sense the rolw of AndroidKeystoreAesGcm. So my understanding is that tink does the following

Tink creates a master key with the alias passed in withMasterKeyUri(). The key is stored in the AndroidKeyStore.
Tink gets an Aead from AndroidKeystoreKmsClient#getOrGenerateNewAeadKey()
Tink writes the KeySet into the SharedPreference with the master key (stored in the KeyStore). This encryption will use AndroidKeystoreAesGcm to encrypt, as it is passed in as the masterKey in KeysetHandle#encrypt(Keyset keyset, Aead masterKey)

Any update on this? Still happens in 1.0.0-rc03 (tested in Mi A2 Lite)

authors: AllanHasegawa GuilhE daewin simtse thaidn tholenst

repo: google/tink


text:
Make sure we do not alpha rename this and arguments in methods

Fixes #1631

LGTM

authors: arv johnjbarton

repo: google/traceur-compiler


text:
Use instead of restoring side table for Loop

On my Linux Docker container:

CoreMark results of this PR: 37.991516 (in 18.695s), 36.7287 (in 19.258s)
CoreMark results of https://github.com/google/wasefire/pull/690: 37.101162 (in 19.133s), 37.39716 (in 18.996s)
Look very similar. Not sure about nordic.

CoreMark results of main: 28.791477 (in 18.09s), 28.814291 (in 17.734s)
So there does seem some minor performance improvement with delta_ip and delta_stp.
https://github.com/google/wasefire/issues/46

On nordic I get:

This pull request (with the suggested changes): 0.13723445 (in 146.037s) and 6660 bytes of RAM
This development branch: 0.13704263 (in 146.247s) and 7588 of RAM
The main branch: 0.09060761 (in 221.04s) and 5416 of RAM

So we should merge this PR, not only because it uses less RAM, but also because it simplifies the code.

How do you measure RAM?

repo: google/wasefire


text:
Drop target action should be a single undo/redo history item

Bug Description
Drop targets are done on two steps: Changing the element's content then deleting the original dropped element. This needs to be combined into one action for undo/redo purposes.
Possible duplicate See #793
Expected Behaviour
Undoing a drop should restore the masked element and also restore the original element. Right now it restores the original element but doesn't restore the masked element's contents until a second undo is issued
Steps to Reproduce

Create a star shape
Add an image to the canvas
Drag the image into the star's drop target
The star should now show the image as its contents
Press the "undo" button
The image comes back into the canvas but the star shape still has the image as its contents

Do not alter or remove anything below. The following sections will be managed by moderators only.
Acceptance Criteria
QA Instructions

Removing from the board, no longer required for alpha

Set as p1 per the above

authors: jauyong wassgha

repo: google/web-stories-wp


text:
Bug: Persist Delay and Duration between Effect Chooser Direction Changes

Summary
Now if you maintain the same animation, but change its direction through the effect chooser, it should maintain duration and delay.
Relevant Technical Choices
NA
To-do
NA
User-facing changes
Duration and delay maintained if you change direction through animation effect chooser
Testing Instructions
-> Go to story editor
-> Select foreground element
-> Apply fly in from the left animation
-> update duration and delay to a new value
-> reopen effect chooser and select fly in but from another direction
-> exit out of the effect chooser and see that the duration & delay values didn't get reset
-> open the effect chooser and select a completely different animation like fade in
-> close the effect chooser and see that delay & duration are set to the recommended values for fade in

Fixes #5523

Verified in QA

authors: csossi littlemilkstudio