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/args hatrun java matmul 

## Run with the CUDA Backend
java @bldr/args hatrun 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

Error

 ⚠️ OCA signatory status must be verified

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

[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]

❗ This change is not yet ready to be integrated.
See the Progress checklist in the description for automated requirements.

[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!