Skip to content

Commit bb04554

Browse files
MikeL270MikeL270
committed
fix onemath_gemm to actually teach gemm and remove buffer/accessor
1 parent 2ce7ec7 commit bb04554

1 file changed

Lines changed: 223 additions & 138 deletions

File tree

‎Lesson_Materials/oneMath_gemm/index.html‎

Lines changed: 223 additions & 138 deletions
Original file line numberDiff line numberDiff line change
@@ -1,165 +1,250 @@
11
<!doctype html>
22

33
<html>
4-
<head>
5-
<meta charset="utf-8"/>
6-
<link rel="stylesheet" href="../../Static/css/reveal.css"/>
7-
<link rel="stylesheet" href="../../Static/css/theme/white.css"/>
8-
<link rel="stylesheet" href="../../Static/css/custom.css"/>
9-
<link rel="stylesheet" href="../../Static/css/atom-one-light.css"/>
10-
<script src="../../node_modules/reveal.js/dist/reveal.js" defer></script>
11-
<script src="../../node_modules/reveal.js/dist/plugin/markdown.js" defer></script>
12-
<script src="../../node_modules/reveal.js/dist/plugin/highlight.js" defer></script>
13-
<script src="../../node_modules/reveal.js/dist/plugin/notes.js" defer></script>
14-
<script src="../../js/plugins/merge-html.js" defer></script>
15-
<script src="../../js/setup.js" defer></script>
16-
</head>
17-
<body>
18-
<div class="reveal">
19-
<div class="slides">
20-
<!--Slide 1-->
21-
<section class="hbox">
22-
<h2 style="text-transform: none;">
23-
oneAPI Math Library (oneMath)
24-
</h2>
25-
</section>
26-
<!--Slide 2-->
27-
<section class="hbox" data-markdown>
28-
## Learning Objectives
29-
* Learn what the oneMath is and how it works
30-
* Learn how to use GEMM APIs from oneMath with both USM and buffer memory models
31-
</section>
32-
<!--Slide 3-->
33-
<section>
34-
<div class="hbox" data-markdown>
35-
## Do you need to write your own kernels?
36-
</div>
374

38-
<div class="container" data-markdown>
39-
* Many computationally intensive applications spend the most of their time in **common operations / algorithms**
5+
<head>
6+
<meta charset="utf-8" />
7+
<link rel="stylesheet" href="../../Static/css/reveal.css" />
8+
<link rel="stylesheet" href="../../Static/css/theme/white.css" />
9+
<link rel="stylesheet" href="../../Static/css/custom.css" />
10+
<link rel="stylesheet" href="../../Static/css/atom-one-light.css" />
11+
<script src="../../node_modules/reveal.js/dist/reveal.js" defer></script>
12+
<script src="../../node_modules/reveal.js/dist/plugin/markdown.js" defer></script>
13+
<script src="../../node_modules/reveal.js/dist/plugin/highlight.js" defer></script>
14+
<script src="../../node_modules/reveal.js/dist/plugin/notes.js" defer></script>
15+
<script src="../../js/plugins/merge-html.js" defer></script>
16+
<script src="../../js/setup.js" defer></script>
17+
</head>
18+
19+
<body>
20+
<div class="reveal">
21+
<div class="slides">
22+
<!--Slide 1-->
23+
<section class="hbox">
24+
<h2 style="text-transform: none;">
25+
oneAPI Math Library (oneMath)
26+
</h2>
27+
</section>
28+
<!--Slide 2-->
29+
<section class="hbox" data-markdown>
30+
## Learning Objectives
31+
* Learn what the oneMath is and how it works
32+
* Learn how oneMath dispatches work to a backend at run time or compile time
33+
* Learn what GEMM is
34+
* Learn how to use the GEMM API from oneMath with the USM memory model
35+
</section>
36+
<!--Slide 3-->
37+
<section>
38+
<div class="hbox" data-markdown>
39+
## Do you need to write your own kernels?
40+
</div>
41+
<div class="container" data-markdown>
42+
* Many computationally intensive applications spend the most of their time in **common operations /
43+
algorithms**
4044
* **Numerical libraries** provide reliable solutions to these common problems
41-
* You can focus on solving higher-level problems instead of technical details
45+
* You can focus on solving higher-level problems instead of technical details
4246
* Libraries optimised for specific hardware provide **superior performance**
43-
</div>
44-
</section>
45-
<!--Slide 4-->
46-
<section>
47-
<div class="hbox" data-markdown>
48-
## Numerical libraries
49-
</div>
50-
51-
<div class="container" data-markdown>
52-
* Common APIs like BLAS or LAPACK have multiple CPU implementations and vendor-specific GPU solutions
53-
* **Intel CPU/GPU**: Intel Math Kernels Library (oneMKL)
54-
* **NVIDIA GPU**: cuBLAS, cuSOLVER, cuRAND, cuFFT
55-
* **AMD GPU**: rocBLAS, rocSOLVER, rocRAND, rocFFT
47+
</div>
48+
</section>
49+
<!--Slide 4-->
50+
<section>
51+
<div class="hbox" data-markdown>
52+
## Numerical libraries
53+
</div>
54+
<div class="container" data-markdown>
55+
* Common APIs like BLAS or LAPACK have multiple CPU implementations and vendor-specific GPU
56+
solutions
57+
* **Intel CPU/GPU**: Intel Math Kernels Library (oneMKL)
58+
* **NVIDIA GPU**: cuBLAS, cuSOLVER, cuRAND, cuFFT
59+
* **AMD GPU**: rocBLAS, rocSOLVER, rocRAND, rocFFT
5660
* Imagine being able to use all of them with *single source code* &#8594; **oneMath**
57-
</div>
58-
</section>
59-
<!--Slide 5-->
60-
<section>
61-
<div class="hbox">
62-
<h2 style="text-transform: none;">
63-
oneAPI and oneMath
64-
</h2>
65-
</div>
66-
<div class="container" data-markdown>
67-
* Open-source [**oneAPI**](https://oneapi.io/) project governed by the [United Acceleration (UXL) Foundation](https://uxlfoundation.org/):
68-
* defines SYCL-based APIs and provides library implementations
69-
* brings performance and ease of development to SYCL applications
70-
* [**oneMath** specification](https://oneapi-spec.uxlfoundation.org/specifications/oneapi/latest/elements/onemath/source/):
71-
* defines SYCL API for numerical computations across several domains
72-
* Linear Algebra, Discrete Fourier Transforms, Random Number Generators, Statistics, Vector Math
73-
* [**oneMath** library](https://github.com/uxlfoundation/oneMath):
74-
* wrapper implementation dispatching SYCL API calls to a multitude of implementations, both generic and vendor-specific
75-
</div>
76-
<div class="container">
77-
<img style="height:4em; margin-top:2em;" src="../../Static/images/oneAPI.png"/>
78-
<div style="display: inline; width: 2em;"></div>
79-
<img style="height:4em; margin-top:2em;" src="../../Static/images/uxl.svg"/>
80-
</div>
81-
</section>
82-
<!--Slide 6-->
83-
<section>
61+
</div>
62+
</section>
63+
<!--Slide 5-->
64+
<section>
65+
<div class="hbox">
8466
<h2 style="text-transform: none;">
85-
oneMath library backends
67+
oneAPI and oneMath
8668
</h2>
87-
<object class="r-stretch" data="../../Static/images/oneMath-backends.svg"></object>
88-
</section>
89-
<!--Slide 7-->
90-
<section>
91-
<div class="hbox" data-markdown>
92-
#### Run-time dispatching
93-
</div>
94-
<div class="container">
95-
<pre><code class="language-cpp" data-line-numbers>
69+
</div>
70+
<div class="container" data-markdown>
71+
* Open-source [**oneAPI**](https://oneapi.io/) project governed by the [United Acceleration (UXL)
72+
Foundation](https://uxlfoundation.org/):
73+
* defines SYCL-based APIs and provides library implementations
74+
* brings performance and ease of development to SYCL applications
75+
* [**oneMath**
76+
specification](https://oneapi-spec.uxlfoundation.org/specifications/oneapi/latest/elements/onemath/source/):
77+
* defines SYCL API for numerical computations across several domains
78+
* Linear Algebra, Discrete Fourier Transforms, Random Number Generators, Statistics, Vector Math
79+
* [**oneMath** library](https://github.com/uxlfoundation/oneMath):
80+
* wrapper implementation dispatching SYCL API calls to a multitude of implementations, both generic
81+
and vendor-specific
82+
</div>
83+
<div class="container">
84+
<img style="height:4em; margin-top:2em;" src="../../Static/images/oneAPI.png" />
85+
<div style="display: inline; width: 2em;"></div>
86+
<img style="height:4em; margin-top:2em;" src="../../Static/images/uxl.svg" />
87+
</div>
88+
</section>
89+
<!--Slide 6-->
90+
<section>
91+
<h2 style="text-transform: none;">
92+
oneMath library backends
93+
</h2>
94+
<object class="r-stretch" data="../../Static/images/oneMath-backends.svg"></object>
95+
</section>
96+
<!--Slide 7-->
97+
<section>
98+
<div class="hbox" data-markdown>
99+
#### Dispatching
100+
</div>
101+
<div class="container">
102+
<div class="col-left-3">
103+
<pre><code class="language-cpp" data-line-numbers="|8,11-14|16-20" data-fragment-index="0" data-noescape data-trim>
104+
#include &lt;cassert&gt;
96105
#include &lt;oneapi/math.hpp&gt;
106+
#include &lt;sycl/sycl.hpp&gt;
97107

98-
sycl::queue q{myDeviceSelector};
108+
int main() {
109+
int n = 16;
110+
float alpha = 2;
111+
<mark class="fragment" data-fragment-index="0">sycl::queue q</mark>;
112+
float *x = sycl::malloc_device&lt;float&gt;(n, q);
99113

100-
sycl::buffer&lt;T,1&gt; a{a_host, m*k};
101-
sycl::buffer&lt;T,1&gt; b{b_host, k*n};
102-
sycl::buffer&lt;T,1&gt; c{c_host, m*n};
114+
// Run-time dispatch: backend chosen from the queue's device
115+
q.fill(x, 1.0f, n).wait();
116+
oneapi::math::blas::column_major::scal(
117+
<mark class="fragment" data-fragment-index="0">q</mark>, n, alpha, x, 1).wait();
103118

104-
// Compute C = A*B+C on the device
105-
oneapi::math::blas::column_major::gemm(q, ..., m, n, k, ..., a, ..., b, ..., c, ... );
119+
// Compile-time dispatch: backend fixed at compile time
120+
oneapi::math::backend_selector&lt;oneapi::math::backend::generic&gt; <mark class="fragment" data-fragment-index="1">sel</mark>{q};
121+
q.fill(x, 1.0f, n).wait();
122+
oneapi::math::blas::column_major::scal(
123+
<mark class="fragment" data-fragment-index="1">sel</mark>, n, alpha, x, 1).wait();
124+
125+
float result;
126+
q.memcpy(&amp;result, x, sizeof(float)).wait();
127+
// x started as 1 and was scaled by alpha
128+
assert(result == alpha);
129+
sycl::free(x, q);
130+
}
106131
</code></pre>
107-
</div>
108-
<div class="container" data-markdown>
109-
* Backend is loaded at run time based on the device associated with the SYCL queue
110-
* Both buffer and USM APIs available (mind the different synchronisation)
111-
* The same binary can run on different hardware with a generic device selector
112-
* Can run on CPU or different GPUs without recompiling
113-
* Link the application with the top-level runtime library: `-lonemath`
114132
</div>
115-
</section>
116-
<!--Slide 8-->
117-
<section>
118-
<div class="hbox" data-markdown>
119-
#### Compile-time dispatching
120-
</div>
121-
<div class="container">
122-
<pre><code class="language-cpp" data-line-numbers>
133+
<div class="col-right-2 stack-fragments">
134+
<div class="fragment fade-out" data-fragment-index="0" data-markdown>
135+
* oneMath chooses which backend runs a routine
136+
* The first argument decides how that backend is selected
137+
</div>
138+
<div class="fragment current-visible" data-fragment-index="0" data-markdown>
139+
* **Run-time dispatch:** pass the `queue` directly as the first argument
140+
* The backend is loaded at run time from the device the queue targets
141+
* The same binary can run on different hardware without recompiling
142+
* Link the application with the top-level runtime library: `-lonemath`
143+
</div>
144+
<div class="fragment current-visible" data-fragment-index="1" data-markdown>
145+
* **Compile-time dispatch:** wrap the queue in a `backend_selector` and pass that instead
146+
* The backend is fixed at compile time
147+
* Reduces the small dispatching overhead at the cost of removed portability
148+
* Link the application with the specific backend library: `-lonemath_blas_mklcpu`
149+
</div>
150+
</div>
151+
</div>
152+
</section>
153+
<!--Slide 8-->
154+
<section>
155+
<div class="hbox" data-markdown>
156+
#### What is GEMM?
157+
</div>
158+
<div class="container" data-markdown>
159+
* GEMM stands for **GE**neral **M**atrix **M**ultiply
160+
* It is the core Level-3 routine of BLAS (Basic Linear Algebra Subprograms)
161+
* It computes `C = alpha * A * B + beta * C`
162+
* `A`, `B` and `C` are dense matrices; `alpha` and `beta` are scalars
163+
* `A` and/or `B` may optionally be transposed
164+
* It is the building block of dense linear algebra and machine learning, so every backend optimises it heavily
165+
</div>
166+
</section>
167+
<!--Slide 9-->
168+
<section>
169+
<div class="hbox" data-markdown>
170+
#### GEMM with oneMath
171+
</div>
172+
<div class="container">
173+
<pre><code class="language-cpp" data-line-numbers="|7-12|16-22|25-27|30-31" data-fragment-index="0" data-noescape data-trim>
174+
#include &lt;cassert&gt;
123175
#include &lt;oneapi/math.hpp&gt;
176+
#include &lt;sycl/sycl.hpp&gt;
177+
#include &lt;vector&gt;
124178

125-
sycl::queue cpu_queue{sycl::cpu_selector_v};
179+
int main() {
180+
// C(m x n) = alpha * A(m x k) * B(k x n) + beta * C(m x n)
181+
int m = 16, n = 16, k = 16;
182+
int lda = k, ldb = n, ldc = n;
183+
float alpha = 1, beta = 0;
184+
auto transA = oneapi::math::transpose::nontrans;
185+
auto transB = oneapi::math::transpose::nontrans;
126186

127-
sycl::buffer&lt;T,1&gt; a{a_host, m*k};
128-
sycl::buffer&lt;T,1&gt; b{b_host, k*n};
129-
sycl::buffer&lt;T,1&gt; c{c_host, m*n};
187+
sycl::queue q;
130188

131-
oneapi::math::backend_selector&lt;oneapi::math::backend::mklcpu&gt; cpu_selector(cpu_queue);
132-
// Select the Intel oneMKL CPU backend specifically ^^^^^^
189+
// Allocate device memory and copy the inputs across
190+
float *dA = sycl::malloc_device&lt;float&gt;(m * k, q);
191+
float *dB = sycl::malloc_device&lt;float&gt;(k * n, q);
192+
float *dC = sycl::malloc_device&lt;float&gt;(m * n, q);
193+
std::vector&lt;float&gt; A(m * k, 1), B(k * n, 1), C(m * n);
194+
q.memcpy(dA, A.data(), sizeof(float) * m * k);
195+
q.memcpy(dB, B.data(), sizeof(float) * k * n).wait();
133196

134-
oneapi::math::blas::column_major::gemm(cpu_selector, ..., m, n, k, ..., a, ..., b, ..., c, ... );
135-
</code></pre>
197+
// Compute C = alpha * A * B + beta * C on the device
198+
oneapi::math::blas::row_major::gemm(
199+
q, transA, transB, m, n, k, alpha,
200+
dA, lda, dB, ldb, beta, dC, ldc).wait();
201+
202+
// Copy the result back and check it
203+
q.memcpy(C.data(), dC, sizeof(float) * m * n).wait();
204+
assert(C[0] == static_cast&lt;float&gt;(k));
205+
206+
sycl::free(dA, q); sycl::free(dB, q); sycl::free(dC, q);
207+
}
208+
</code></pre>
209+
</div>
210+
<div class="bottom-bullets stack-fragments">
211+
<div class="fragment fade-out" data-fragment-index="0" data-markdown>
212+
* oneMath exposes GEMM as a single call, with the matrices held in USM device memory
136213
</div>
137-
<div class="container" data-markdown>
138-
* Specific backend can be selected at compile-time with a `backend_selector`
139-
* Passed into the API in place of the queue
140-
* Reduces the small dispatching overhead at the cost of removed portability
141-
* Link the application with the specific backend library: `-lonemath_blas_mklcpu`
142-
</div>
143-
</section>
144-
<!--Slide 9-->
145-
<section>
146-
<div class="hbox" data-markdown>
147-
## Exercise
214+
<div class="fragment current-visible" data-fragment-index="0" data-markdown>
215+
* The GEMM call is described by the dimensions `m`, `n`, `k`, the leading dimensions, the `alpha`/`beta` scalars and whether `A` or `B` are transposed
148216
</div>
149-
<div class="container" data-markdown>
150-
* Objectives: Learn to use oneMath GEMM buffer and USM APIs
217+
<div class="fragment current-visible" data-fragment-index="1" data-markdown>
218+
* Allocate device memory with USM and copy the input matrices across
219+
</div>
220+
<div class="fragment current-visible" data-fragment-index="2" data-markdown>
221+
* A single call computes `C = alpha * A * B + beta * C`; `row_major` matches C/C++ array layout
222+
</div>
223+
<div class="fragment current-visible" data-fragment-index="3" data-markdown>
224+
* Copy the result back and check it: `A` and `B` are all ones, so every element of `C` is `k`
225+
</div>
226+
</div>
227+
</section>
228+
<!--Slide 10-->
229+
<section>
230+
<div class="hbox" data-markdown>
231+
## Exercise
232+
</div>
233+
<div class="container" data-markdown>
234+
* Objectives: Learn to use the oneMath GEMM USM API
151235
* Boiler-plate code already provided to:
152-
* Initialize matrices on host
153-
* Compute reference result on host
154-
* Compare the host and device results
236+
* Initialize matrices on host
237+
* Compute reference result on host
238+
* Compare the host and device results
155239
* Please **complete the TODO tasks** marked in the `source_*.cpp`
156-
* Create buffers or transfer data with USM
157-
* Compute GEMM by calling the oneMath API
158-
* Use the provided `VerifyResult` function
240+
* Transfer data to the device with USM
241+
* Compute GEMM by calling the oneMath API
242+
* Use the provided `VerifyResult` function
159243
* If stuck, have a look at `solution_*.cpp`
160-
</div>
161-
</section>
162-
</div>
244+
</div>
245+
</section>
163246
</div>
164-
</body>
247+
</div>
248+
</body>
249+
165250
</html>

0 commit comments

Comments
 (0)