alexander.nutz/website
1
0
Fork 0
mirror of https://github.com/alex-s168/website.git synced 2025-09-10 01:05:07 +02:00
Code Issues Packages Projects Releases Wiki Activity

c

Browse Source
This commit is contained in:
Alexander Nutz
2025-08-26 21:13:01 +02:00
parent 89c9b3069e
commit 5a9dfdd720
2 changed files with 16 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
pages/article-gpu-arch-1.typ
View File

@@ -280,6 +280,8 @@
{ {
int row = get_global_id(1); // y int row = get_global_id(1); // y
int col = get_global_id(0); // x int col = get_global_id(0); // x
int local_row = get_local_id(1); // y
int local_col = get_local_id(0); // x
__local float Asub[TILE_SIZE][TILE_SIZE]; __local float Asub[TILE_SIZE][TILE_SIZE];
__local float Bsub[TILE_SIZE][TILE_SIZE]; __local float Bsub[TILE_SIZE][TILE_SIZE];
@@ -289,24 +291,28 @@
for (int t = 0; t < N / TILE_SIZE; ++t) { for (int t = 0; t < N / TILE_SIZE; ++t) {
// load tiles into local // load tiles into local
int tiledRow = row; int tiledRow = row;
int tiledCol = t * TILE_SIZE + get_local_id(0); int tiledCol = t * TILE_SIZE + local_col;
float av;
if (tiledRow < N && tiledCol < N) if (tiledRow < N && tiledCol < N)
Asub[get_local_id(1)][get_local_id(0)] = A[tiledRow * N + tiledCol]; av = A[tiledRow * N + tiledCol];
else else
Asub[get_local_id(1)][get_local_id(0)] = 0.0f; av = 0.0f;
Asub[local_row][local_col] = av;
tiledRow = t * TILE_SIZE + get_local_id(1); tiledRow = t * TILE_SIZE + local_row;
tiledCol = col; tiledCol = col;
float bv;
if (tiledRow < N && tiledCol < N) if (tiledRow < N && tiledCol < N)
Bsub[get_local_id(1)][get_local_id(0)] = B[tiledRow * N + tiledCol]; bv; = B[tiledRow * N + tiledCol];
else else
Bsub[get_local_id(1)][get_local_id(0)] = 0.0f; bv = 0.0f;
Bsub[local_row][local_col]= bv;
// sync local access across local grp // sync local access across local grp
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for (int k = 0; k < TILE_SIZE; ++k) for (int k = 0; k < TILE_SIZE; ++k)
sum += Asub[get_local_id(1)][k] * Bsub[k][get_local_id(0)]; sum += Asub[local_row][k] * Bsub[k][local_col];
// sync local access across local grp // sync local access across local grp
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);

4
pages/index.typ
View File

@@ -87,9 +87,11 @@
#link("https://github.com/alex-s168/website")[Website source code] #link("https://github.com/alex-s168/website")[Website source code]
#br()#br() #br()#br()
The latest version of my badge will always be at:\ Latest version of my badge:\
#raw(people.alex.badge) #raw(people.alex.badge)
#br()
Check out these websites:
#context if is-web and is-html() { // excludes min.html builds too #context if is-web and is-html() { // excludes min.html builds too
let scale = 1.3 let scale = 1.3
for id in people.keys() { for id in people.keys() {