| | |
| | | |
| | | float val = 0; |
| | | |
| | | int row_block = get_group_id(0); |
| | | int col_block = get_group_id(1); |
| | | int row_block = get_group_id(1); |
| | | int col_block = get_group_id(0); |
| | | |
| | | int sub_row = get_local_id(0); |
| | | int sub_col = get_local_id(1); |
| | | int sub_row = get_local_id(1); |
| | | int sub_col = get_local_id(0); |
| | | |
| | | int row = row_block*BLOCK + sub_row; |
| | | int col = col_block*BLOCK + sub_col; |
| | |
| | | int brow = i + sub_row; |
| | | int bcol = col_block*BLOCK + sub_col; |
| | | |
| | | Asub[sub_row][sub_col] = TA ? A[arow + acol*lda] : A[arow*lda + acol]; |
| | | Bsub[sub_row][sub_col] = TB ? B[brow + bcol*ldb] : B[brow*ldb + bcol]; |
| | | if(arow < M && acol < K)Asub[sub_row][sub_col] = TA ? A[arow + acol*lda] : A[arow*lda + acol]; |
| | | if(brow < K && bcol < N)Bsub[sub_row][sub_col] = TB ? B[brow + bcol*ldb] : B[brow*ldb + bcol]; |
| | | |
| | | barrier(CLK_LOCAL_MEM_FENCE); |
| | | |