130 lines
3.7 KiB
Common Lisp
130 lines
3.7 KiB
Common Lisp
// super super slow sgemm kernel by heehoon
|
|
#define TS 32
|
|
#define WPTF 8
|
|
#define RTSF (TS/WPTF)
|
|
#define WPTL 16
|
|
#define RTSL (TS/WPTL)
|
|
#define PADDINGX 16
|
|
#define PADDINGY 16
|
|
|
|
__kernel void sgemm(__global float *A, __global float *B, __global float *C, int M, int N, int K) {
|
|
const int row = get_local_id(0); // row index of C
|
|
const int col = get_local_id(1); // row index of C
|
|
const int global_row = TS*get_group_id(0)+row;
|
|
const int global_col = TS*get_group_id(1)+col;
|
|
__local float Asub[TS][TS];
|
|
__local float Bsub[TS][TS];
|
|
|
|
float acc[WPTF];
|
|
for(int w=0; w<WPTF; w++) {
|
|
acc[w] = 0.0f;
|
|
}
|
|
|
|
const int num_tiles = (K%TS)>0 ? K/TS+1 : K/TS;
|
|
|
|
for(int t=0;t<num_tiles;t++){
|
|
for(int w=0;w<WPTF;w++){
|
|
const int tile_row = TS*t+row;
|
|
const int tile_col = TS*t+col;
|
|
|
|
if(global_row+w*RTSF>=M || tile_col >= K)
|
|
Asub[row+w*RTSF][col]=0.0f;
|
|
else
|
|
Asub[row+w*RTSF][col]=A[(global_row+w*RTSF)*K+tile_col];
|
|
|
|
if(tile_row+w*RTSF>=K||global_col>=N)
|
|
Bsub[row+w*RTSF][col]=0.0f;
|
|
else
|
|
Bsub[row+w*RTSF][col]=B[(tile_row+w*RTSF)*N+global_col];
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
for(int k = 0; k < TS; k++) {
|
|
for(int w=0;w<WPTF;w++) {
|
|
acc[w] += Asub[row+w*RTSF][k]*Bsub[k][col];
|
|
}
|
|
}
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
for(int w=0; w<WPTF; w++) {
|
|
if(global_row+w*RTSF>=M || global_col >=N) continue;
|
|
else C[(global_row+w*RTSF)*N+global_col]=acc[w];
|
|
}
|
|
}
|
|
|
|
__kernel void sgemma(__global float *A, __global float *B, __global float *C, int M, int N, int K) {
|
|
const int row = get_local_id(0); // row index of C
|
|
const int col = get_local_id(1); // row index of C
|
|
const int global_row = TS*get_group_id(0)+row;
|
|
const int global_col = TS*get_group_id(1)+col;
|
|
__local float Asub[TS][TS];
|
|
__local float Bsub[TS][TS];
|
|
|
|
float acc[WPTL];
|
|
for(int w=0;w<WPTL;w++) {
|
|
acc[w] = 0.0f;
|
|
}
|
|
|
|
const int num_tiles = K/TS;
|
|
|
|
for(int t=0;t<num_tiles;t++){
|
|
for(int w=0;w<WPTL;w++){
|
|
const int tile_row = TS*t+row;
|
|
const int tile_col = TS*t+col;
|
|
|
|
Asub[row+w*RTSL][col]=A[(global_row+w*RTSL)*K+tile_col];
|
|
Bsub[row+w*RTSL][col]=B[(tile_row+w*RTSL)*N+global_col];
|
|
}
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
|
|
for(int k = 0; k < TS; k++) {
|
|
for(int w=0;w<WPTL;w++) {
|
|
acc[w] += Asub[row+w*RTSL][k]*Bsub[k][col];
|
|
}
|
|
}
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE);
|
|
}
|
|
|
|
for(int w=0;w<WPTL;w++) {
|
|
C[(global_row+w*RTSL)*N+global_col]=acc[w];
|
|
}
|
|
}
|
|
|
|
__kernel void paddingAddZeroes(const int P, const int Q,
|
|
const __global float* input,
|
|
const int P_XL, const int Q_XL,
|
|
__global float* output) {
|
|
|
|
const int tx = get_group_id(0)*PADDINGX + get_local_id(0); // 0..P_XL in blocks of PADDINGX
|
|
const int ty = get_group_id(1)*PADDINGY + get_local_id(1); // 0..Q_XL in blocks of PADDINGY
|
|
|
|
if (tx < P_XL && ty < Q_XL) {
|
|
float value;
|
|
if (tx < P && ty < Q) {
|
|
value = input[ty*P + tx];
|
|
}
|
|
else {
|
|
value = 0.0f;
|
|
}
|
|
|
|
output[ty*P_XL + tx] = value;
|
|
}
|
|
}
|
|
|
|
__kernel void paddingRemoveZeroes(const int P_XL, const int Q_XL,
|
|
const __global float* input,
|
|
const int P, const int Q,
|
|
__global float* output) {
|
|
|
|
const int tx = get_group_id(0)*PADDINGX + get_local_id(0); // 0..P in blocks of PADDINGX
|
|
const int ty = get_group_id(1)*PADDINGY + get_local_id(1); // 0..Q in blocks of PADDINGY
|
|
|
|
if (tx < P && ty < Q) {
|
|
output[ty*P + tx] = input[ty*P_XL + tx];
|
|
}
|
|
}
|
|
|