GV

GeezerVison

The World Wide Web, restored to 800 pixels of glory.

Home|About|Gaming cabinet|Archives|How it works|Starter sites

On this site

  • Sign in
  • ActionsCI/CD and automation
  • CodespacesCloud development environments
  • All Features
  • Documentation
  • Training
  • Certifications
  • Build with Copilot
  • Customize Copilot
  • Integrations
  • Trending repos
  • GitHub Sponsors
  • Developer articles
  • Enterprises

This view

  • Original page
  • New address

GeezerVison is showing github.com — Fixes to mean and better precision on sum (cpu) by willyborn · Pull Request #3687 · arrayfire/arrayfire

Conversation

@willyborn

@willyborn willyborn commented Aug 4, 2025

Copy link
Copy Markdown
Contributor

When strided weights are used, the mean function typically returned random results.
The produced result of mean is no longer dependent on platform or processing style (linear, strided, on CPU or GPU)

Description

b74f28f sum on CPU now has the same precision as other platforms

  • Introduction of Kahan summation algorithm for series. This algorithm reduced the error when a series of summations are performed with values of varying sizes.
    With this introduction, the difference in result of a linear array and strided array is minimal. The precision in the tests are as a result improved.
    As result, the precision of many tests can be improved now since all platforms generate similar results, independent from parallel or serial processing.

9307d6f Fixes to mean, on all platforms
All:

  • Introduction of Kahan summation algorithm for series. This algorithm reduced the error when a series of summations are performed with values of varying sizes.
    With this introduction, the difference in result of a linear array and strided array is minimal. The precision in the tests are as a result improved.
  • Negative weights can no longer generate a division by zero.
  • In most cases sub-arrays would generated random results with strided weights.

CPU:

  • the idxs of the weights array is now calculated according to its own strides iso the strides of the input array.

CUDA:

  • Param were used without full initialisation. Because the construction of Param’s are deprecated, they are replaced by the createEmptyArray() constructions.
  • Idxs of weights are now calculated according to its own strides. Idem for the offsets.
  • Small arrays are processed on the CPU. Strides arrays are now processed on the GPU, independent from size since the CPU version only handles linear arrays.

ONEAPI:

  • Idxs of weights are now calculated according to its own strides. Idem for offsets.
  • Param were used without full initialisation. Because the construction of Param’s are deprecated, they are replaced by the createEmptyArray() constructions.
  • Small arrays are processed on the CPU. Strides arrays are now processed on the GPU, independent from size since the CPU version only handles linear arrays.
  • Offsets are handle for the linear processing.

OPENCL:

  • Idxs of weights are now calculated according to its own strides. Idem for offsets.
  • When converting a linear array to a vector, the weights offsets were overwritten.
  • The size of the tmpOut and tmpWeight arrays are now correctly sized. Intermediate results were overwritten resulting in a mean of a part of the array.
  • Small arrays are processed on the CPU. Strides arrays are now processed on the GPU, independent from size since the CPU version only handles linear arrays.
  • Param were used without full initialisation. Because the construction of Param’s are deprecated, they are replaced by the createEmptyArray() constructions.

TEST:
- extra tests added on all temporary formats.
- allowed fault tolerance for tests is lowered, since now the correct mean is calculated for all backends.
- MeanAllTest now uses random data iso constants for testing. This blocked the detection of partial processing of the input.

Additional information about the PR answering following questions:

  • Is this a new feature or a bug fix? BUG
  • More detail if necessary to describe all commits in pull request.
  • Why these changes are necessary. CORRECT RESULTS WERE NOT GUARANTEED
  • Potential impact on specific hardware, software or backends. NONE
  • New functions and their functionality. NONE
  • Can this PR be backported to older versions? NO
  • Future changes not implemented in this PR. NO
    Fixes:

Changes to Users

Better precision in all circumstances, and independent form platform or array format.

Checklist

  • Rebased on latest master
  • Code compiles
  • Tests pass
  • Functions added to unified API
  • Functions documented

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

All reactions
@christophe-murphy

christophe-murphy commented Aug 4, 2025

Copy link
Copy Markdown
Contributor

Thanks for this willyborn. FYI, our Buildbot CI system failed to build the PR with the following error:

/var/lib/buildbot/workers/ubuntu2204-cuda-toolkit-12-9/linux-all/build/src/backend/cuda/kernel/mean.hpp(65): error: more than one operator "*" matches these operands: built-in operator "arithmetic * arithmetic" function "operator*(const __half &, const __half &)" (declared at line 223 of /usr/local/cuda/include/cuda_fp16.hpp) operand types are: __half * float To y = (rhs - (*lhs)) * r_wt / (*l_wt) - (*c);

willyborn reacted with thumbs up emoji
All reactions
  • 👍 1 reaction

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

@willyborn
sum on CPU now has the same precision as other platforms
058db8a
@willyborn
willyborn force-pushed the subArray_mean branch from b74f28f to b449cbf Compare August 6, 2025 18:58
@willyborn

willyborn commented Aug 6, 2025

Copy link
Copy Markdown
Contributor Author

I'm on ubuntu 22.04LTS with GCC, and I can not reproduce any warning/error you mentioned.
As result, I encapsulated every element inside the stable_mean so that misunderstandings by the compiler now are eliminated. I hope this solves the error.

PS: Just to be certain I performed the same change for all platforms.

All reactions

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

@christophe-murphy

christophe-murphy commented Aug 7, 2025

Copy link
Copy Markdown
Contributor

This is happening when NVCC is compiling the stable_mean CUDA device kernel. Which version of CUDA are you using? Could you try with 12.9

All reactions

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

@willyborn
Fixes to mean, on all platforms
4a8dca4
@willyborn
willyborn force-pushed the subArray_mean branch from b449cbf to 4a8dca4 Compare August 8, 2025 18:02
@willyborn

willyborn commented Aug 8, 2025 •
edited
Loading

Uh oh!

There was an error while loading. Please reload this page.

Copy link
Copy Markdown
Contributor Author

I am using CUDA 12.9.86 on ubuntu 24.04.3 LTS kernel 6.14.
I have a GTX1080 which is on architecture 6.1 (stores halves, calculates in float).

As soon as I included all major architectures to the nvcc options, I got the error also, which is now also fixed and committed.

All reactions

Sorry, something went wrong.

Uh oh!

There was an error while loading. Please reload this page.

@melonakos melonakos mentioned this pull request Sep 12, 2026
cpu: evaluate arrays past 2^31 elements and sum with compensation #3731
Draft
@melonakos melonakos mentioned this pull request Oct 24, 2026
ci: restore build checks and CPU summation accuracy #3735
Open

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Uh oh!

There was an error while loading. Please reload this page.

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

Uh oh!

There was an error while loading. Please reload this page.

2 participants

@willyborn @christophe-murphy
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
You are visitor 01402503 — GeezerVison does not store the pages you look up.
Best viewed with Netscape Navigator 4.7 or Internet Explorer 6 at 800×600. No cookies required.