-
Notifications
You must be signed in to change notification settings - Fork 32
Expand file tree
/
Copy pathmatrixTranspose_v3.cl
More file actions
42 lines (34 loc) · 1.54 KB
/
Copy pathmatrixTranspose_v3.cl
File metadata and controls
42 lines (34 loc) · 1.54 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
__kernel void matrixTranspose(const int heightA,
const int widthA,
__global const float *a,
__global float *a_T) {
const int alpha = 4;
const int colA = get_global_id(0) * alpha;
const int rowA = get_global_id(1) * alpha;
if ((rowA < heightA) && (colA < widthA)) {
float *src_a_0 = a + (rowA+0) * widthA + (colA+0);
float *src_a_1 = a + (rowA+1) * widthA + (colA+0);
float *src_a_2 = a + (rowA+2) * widthA + (colA+0);
float *src_a_3 = a + (rowA+3) * widthA + (colA+0);
float4 vec1x4_0 = *((float4 *)src_a_0);
float4 vec1x4_1 = *((float4 *)src_a_1);
float4 vec1x4_2 = *((float4 *)src_a_2);
float4 vec1x4_3 = *((float4 *)src_a_3);
a_T[(colA+0) * heightA + (rowA+0)] = vec1x4_0.x;
a_T[(colA+1) * heightA + (rowA+0)] = vec1x4_0.y;
a_T[(colA+2) * heightA + (rowA+0)] = vec1x4_0.z;
a_T[(colA+3) * heightA + (rowA+0)] = vec1x4_0.w;
a_T[(colA+0) * heightA + (rowA+1)] = vec1x4_1.x;
a_T[(colA+1) * heightA + (rowA+1)] = vec1x4_1.y;
a_T[(colA+2) * heightA + (rowA+1)] = vec1x4_1.z;
a_T[(colA+3) * heightA + (rowA+1)] = vec1x4_1.w;
a_T[(colA+0) * heightA + (rowA+2)] = vec1x4_2.x;
a_T[(colA+1) * heightA + (rowA+2)] = vec1x4_2.y;
a_T[(colA+2) * heightA + (rowA+2)] = vec1x4_2.z;
a_T[(colA+3) * heightA + (rowA+2)] = vec1x4_2.w;
a_T[(colA+0) * heightA + (rowA+3)] = vec1x4_3.x;
a_T[(colA+1) * heightA + (rowA+3)] = vec1x4_3.y;
a_T[(colA+2) * heightA + (rowA+3)] = vec1x4_3.z;
a_T[(colA+3) * heightA + (rowA+3)] = vec1x4_3.w;
}
}