Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix Aliasing issue in OpenCL #2943

Merged
merged 11 commits into from
Sep 28, 2023
Merged

Conversation

SteveBronder
Copy link
Collaborator

Summary

Fixes #2942 by adding compound assignment operators to result_cl so that assignments in the generated OpenCL for adjoint calculations use += such as the below

/**
 * A kernel for the equivalent of Matrix += Matrix
 * @param[in/out] var2_global Pointer to array of doubles representing matrix
 * @param[in] var2_rows Rows of var2
 * @param[in] var2_view Whether the matrix is upper or lower triangular or plain
 * @param[in] var3_global Pointer to array of doubles representing matrix
 * @param[in] var3_rows Rows of var3
 * @param[in] var3_view Whether the matrix is upper or lower triangular or plain
 */
kernel void calculate(__global double* var2_global, int var2_rows, int var2_view,
    __global double* var3_global, int var3_rows, int var3_view){
    // Get the global ids for the cell we are working in
    int i = get_global_id(0);
    int j = get_global_id(1);
    // If var2 is lower or upper triangular and this cell is 
    //  not in the range of work, assign it a value of zero
    //  Otherwise assign it the value of var2 the cell works with
    double var2 = 0; 
    if (!((!contains_nonzero(var2_view, LOWER) && j < i) || (!contains_nonzero(var2_view, UPPER) && j > i))) {
        var2 = var2_global[i + var2_rows * j];
    }
    // Same for var3
    double var3 = 0; 
    if (!((!contains_nonzero(var3_view, LOWER) && j < i) || (!contains_nonzero(var3_view, UPPER) && j > i))) {
        var3 = var3_global[i + var3_rows * j];
    }
    // Make a temporary storing value of var2 + var3
    double var1 = var2 + var3;
    // Compound Assign var1 to var2
    var2_global[i + var2_rows * j] += var1;
    double var4 = var2 + var3;
// Compound Assign var4 to var2
    var2_global[i + var2_rows * j] += var4;
}

Tests

Tests are added for the compound assignment and can be run via

python ./runTests.py -j12 ./test/unit/math/opencl/kernel_generator/assignment_ops_test.cpp 

Side Effects

Locally I've been getting spurious errors for some of the reverse mode pdf distributions not matching the cpu versions? Though none of those touch the compound operators so I'm not sure why that is happening

Release notes

Fixes aliasing issue in adjoint acccumulation for OpenCL

Checklist

  • Math issue Aliasing issue in OpenCL assignment_impl #2942

  • Copyright holder: Steve Bronder

    The copyright holder is typically you or your assignee, such as a university or company. By submitting this pull request, the copyright holder is agreeing to the license the submitted work under the following licenses:
    - Code: BSD 3-clause (https://opensource.org/licenses/BSD-3-Clause)
    - Documentation: CC-BY 4.0 (https://creativecommons.org/licenses/by/4.0/)

  • the basic tests are passing

    • unit tests pass (to run, use: ./runTests.py test/unit)
    • header checks pass, (make test-headers)
    • dependencies checks pass, (make test-math-dependencies)
    • docs build, (make doxygen)
    • code passes the built in C++ standards checks (make cpplint)
  • the code is written in idiomatic C++ and changes are documented in the doxygen

  • the new changes are tested

… Creates a assignment op tag that is used by adjoint_results to do a += instead of a = into the adjoint matrix
@SteveBronder SteveBronder changed the title Fix/plusequals assign opencl Fix Aliasing issue in OpenCL Sep 13, 2023
@SteveBronder
Copy link
Collaborator Author

@rok-cesnovar or @t4c1 would one of you have time to look at this?

@rok-cesnovar
Copy link
Member

Sure thing, will have a look in the next days.

@SteveBronder
Copy link
Collaborator Author

@rok-cesnovar can you take a look at this?

Copy link
Member

@rok-cesnovar rok-cesnovar left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks great. I was not getting any issues running the tests locally, so I think this is good to go.

@SteveBronder SteveBronder merged commit be485f4 into develop Sep 28, 2023
8 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Aliasing issue in OpenCL assignment_impl
3 participants