|
111
|
1 /* { dg-do compile } */
|
|
|
2 /* { dg-require-effective-target offload_hsa } */
|
|
|
3 /* { dg-options "-fopenmp -fdump-tree-omplower-details" } */
|
|
|
4
|
|
|
5 #define BLOCK_SIZE 16
|
|
|
6
|
|
|
7
|
|
|
8 void tiled_sgemm_tt(const int M, const int N, const int K, const float alpha, const float*A, const int LDA,
|
|
|
9 const float*B, const int LDB, const float beta, float*C, const int LDC){
|
|
|
10
|
|
|
11 #pragma omp target teams map(to:A[M*K],B[K*N]) map(from:C[M*N])
|
|
|
12 #pragma omp distribute collapse(2)
|
|
|
13 for (int C_row_start=0 ; C_row_start < M ; C_row_start+=BLOCK_SIZE)
|
|
|
14 for (int C_col_start=0 ; C_col_start < N ; C_col_start+=BLOCK_SIZE)
|
|
|
15 {
|
|
|
16 // Each team has a local copy of these mini matrices
|
|
|
17 float As[BLOCK_SIZE][BLOCK_SIZE];
|
|
|
18 float Bs[BLOCK_SIZE][BLOCK_SIZE];
|
|
|
19 #pragma omp parallel
|
|
|
20 {
|
|
|
21 int C_row, C_col;
|
|
|
22 float Cval = 0.0;
|
|
|
23
|
|
|
24 for (int kblock = 0; kblock < K ; kblock += BLOCK_SIZE )
|
|
|
25 {
|
|
|
26 #pragma omp for collapse(2)
|
|
|
27 for (int row=0 ; row < BLOCK_SIZE ; row++)
|
|
|
28 for (int col=0 ; col < BLOCK_SIZE ; col++)
|
|
|
29 {
|
|
|
30 C_row = C_row_start + row;
|
|
|
31 C_col = C_col_start + col;
|
|
|
32 if ((C_row < M) && (kblock + col < K))
|
|
|
33 As[row][col] = A[(C_row*LDA)+ kblock + col];
|
|
|
34 else
|
|
|
35 As[row][col] = 0;
|
|
|
36 if ((kblock + row < K) && C_col < N)
|
|
|
37 Bs[row][col] = B[((kblock+row)*LDB)+ C_col];
|
|
|
38 else
|
|
|
39 Bs[row][col] = 0;
|
|
|
40 }
|
|
|
41
|
|
|
42 #pragma omp for collapse(2)
|
|
|
43 for (int row=0 ; row < BLOCK_SIZE ; row++)
|
|
|
44 for (int col=0 ; col < BLOCK_SIZE ; col++)
|
|
|
45 {
|
|
|
46 for (int e = 0; e < BLOCK_SIZE; ++e)
|
|
|
47 Cval += As[row][e] * Bs[e][col];
|
|
|
48 }
|
|
|
49 } /* End for kblock .. */
|
|
|
50
|
|
|
51
|
|
|
52 #pragma omp for collapse(2)
|
|
|
53 for (int row=0 ; row < BLOCK_SIZE ; row++)
|
|
|
54 for (int col=0 ; col < BLOCK_SIZE ; col++)
|
|
|
55 {
|
|
|
56 C_row = C_row_start + row;
|
|
|
57 C_col = C_col_start + col;
|
|
|
58 if ((C_row < M) && (C_col < N))
|
|
|
59 C[(C_row*LDC)+C_col] = alpha*Cval + beta*C[(C_row*LDC)+C_col];
|
|
|
60
|
|
|
61 }
|
|
|
62 } /* end parallel */
|
|
|
63 } /* end target teams distribute */
|
|
|
64 }
|
|
|
65
|
|
|
66 /* { dg-final { scan-tree-dump "Target construct will be turned into a gridified HSA kernel" "omplower" } } */
|
