1D Matrix Multiplication example for HAT

[jjfumero]

Add new example for 1D Matrix Multiplication in HAT.

How to test?

## Compile 
java --add-modules jdk.incubator.code --enable-preview --source 24 bld

## Run with the OpenCL Backend
java @bldr/hatrun ffi-opencl matmul  

## Run with the CUDA Backend
java @bldr/hatrun ffi-ptx matmul 

Note that the generated kernel for OpenCL contains a race condition:

__kernel void matrixMultiplyKernel(
    __global KernelContext_t *kc, __global F32Array_t* matrixA, __global F32Array_t* matrixB, __global F32Array_t* matrixC, int size
){
    kc->x=get_global_id(0);                   //  << Shared struct across all threads to store the thread-id 
    if(kc->x<kc->maxX){
        for(int j = 0; j<size; j=j+1){
            float acc = (float)0;
            for(int k = 0; k<size; k=k+1){
                acc=acc+matrixA->array[(long)(kc->x*size+k)]*matrixB->array[(long)(k*size+j)];
            }
            matrixC->array[(long)(kc->x*size+j)]=acc;
        }
    }
    return;
}

After applying a patch provided by Gary Frost to solve the race condition, it works.

Patch:

diff --git a/hat/hat/src/main/java/hat/backend/c99codebuilders/C99HatKernelBuilder.java b/hat/hat/src/main/java/hat/backend/c99codebuilders/C99HatKernelBuilder.java
index ade90914d7e..2719fed31ed 100644
--- a/hat/hat/src/main/java/hat/backend/c99codebuilders/C99HatKernelBuilder.java
+++ b/hat/hat/src/main/java/hat/backend/c99codebuilders/C99HatKernelBuilder.java
@@ -26,7 +26,6 @@
 
 
 import hat.buffer.Buffer;
-import hat.buffer.KernelContext;
 import hat.callgraph.KernelCallGraph;
 import hat.callgraph.KernelEntrypoint;
 import hat.optools.FuncOpWrapper;
@@ -72,9 +71,13 @@ T typedefStructOrUnion(boolean isStruct, String name, Consumer<T> consumer) {
 
 
     public final T scope() {
-        return
-                identifier("kc").rarrow().identifier("x").equals().globalId().semicolon().nl();
-                //.identifier("kc").rarrow().identifier("maxX").equals().globalSize().semicolon().nl();
+
+        identifier("KernelContext_t").space().identifier("mine").semicolon().nl();
+        identifier("KernelContext_t").asterisk().space().identifier("kc").equals().ampersand().identifier("mine").semicolon().nl();
+        identifier("kc").rarrow().identifier("x").equals().globalId().semicolon().nl();
+        identifier("kc").rarrow().identifier("maxX").equals().identifier("global_kc").rarrow().identifier("maxX").semicolon().nl();
+        return self();
+
     }
 
     public abstract T globalPtrPrefix();
@@ -137,7 +140,7 @@ public T kernelEntrypoint(KernelEntrypoint kernelEntrypoint, Object[] args) {
                 }
             }
             parenNlIndented(_ -> {
-                        globalPtrPrefix().space().suffix_t("KernelContext").space().asterisk().identifier("kc");
+                        globalPtrPrefix().space().suffix_t("KernelContext").space().asterisk().identifier("global_kc");
                         list.stream().skip(1).forEach(info ->
                                 comma().space().type(info.javaType).space().varName(info.varOp)
                         );

Note: this PR does not provide this path, only the example and the runner extension to run the matrix multiplication.

---

Progress

• Change must not contain extraneous whitespace

Reviewing

Using git

Checkout this PR locally:
$ git fetch https://git.openjdk.org/babylon.git pull/276/head:pull/276
$ git checkout pull/276

Update a local copy of the PR:
$ git checkout pull/276
$ git pull https://git.openjdk.org/babylon.git pull/276/head

Using Skara CLI tools

Checkout this PR locally:
$ git pr checkout 276

View PR using the GUI difftool:
$ git pr show -t 276

Using diff file

Download this PR as a diff file:
https://git.openjdk.org/babylon/pull/276.diff

Using Webrev

Link to Webrev Comment

[bridgekeeper]

Hi @jjfumero, welcome to this OpenJDK project and thanks for contributing!

We do not recognize you as Contributor and need to ensure you have signed the Oracle Contributor Agreement (OCA). If you have not signed the OCA, please follow the instructions. Please fill in your GitHub username in the "Username" field of the application. Once you have signed the OCA, please let us know by writing /signed in a comment in this pull request.

If you already are an OpenJDK Author, Committer or Reviewer, please click here to open a new issue so that we can record that fact. Please use "Add GitHub user jjfumero" as summary for the issue.

If you are contributing this work on behalf of your employer and your employer has signed the OCA, please let us know by writing /covered in a comment in this pull request.

[openjdk]

@jjfumero This change now passes all automated pre-integration checks.

ℹ️ This project also has non-automated pre-integration requirements. Please see the file CONTRIBUTING.md for details.

After integration, the commit message for the final commit will be:

1D Matrix Multiplication example for HAT

You can use pull request commands such as /summary, /contributor and /issue to adjust it as needed.

At the time when this comment was updated there had been no new commits pushed to the code-reflection branch. If another commit should be pushed before you perform the /integrate command, your PR will be automatically rebased. If you prefer to avoid any potential automatic rebasing, please check the documentation for the /integrate command for further details.

As you do not have Committer status in this project an existing Committer must agree to sponsor your change.

➡️ To flag this PR as ready for integration with the above commit message, type /integrate in a new comment. (Afterwards, your sponsor types /sponsor in a new comment to perform the integration).

[jjfumero]

/signed

[bridgekeeper]

Thank you! Please allow for up to two weeks to process your OCA, although it is usually done within one to two business days. Also, please note that pull requests that are pending an OCA check will not usually be evaluated, so your patience is appreciated!

[bridgekeeper]

@jjfumero This pull request has been inactive for more than 4 weeks and will be automatically closed if another 4 weeks passes without any activity. To avoid this, simply add a new comment to the pull request. Feel free to ask for assistance if you need help with progressing this pull request towards integration!

[jjfumero]

Still waiting for the OCA approval.

[bridgekeeper]

@jjfumero This pull request has been inactive for more than 4 weeks and will be automatically closed if another 4 weeks passes without any activity. To avoid this, simply add a new comment to the pull request. Feel free to ask for assistance if you need help with progressing this pull request towards integration!

[jjfumero]

Still waiting for the OCA approval

[SirYwell]

You might want to contact [email protected] regarding your OCA approval status. Not sure what's up there, but it shouldn't take that long.

[mlbridge]

Webrevs

• 01: Full (cd3c7ce9)
• 00: Full (65b647fc)

[bridgekeeper]

@jjfumero This pull request has been inactive for more than 4 weeks and will be automatically closed if another 4 weeks passes without any activity. To avoid this, simply add a new comment to the pull request. Feel free to ask for assistance if you need help with progressing this pull request towards integration!

[jjfumero]

Pending for review