1
取自樣本的代碼。我用它創建了一個項目，它工作，但我不明白一些部分。OpenCL轉置內核如何使用get_local_id
對於該示例的目的，說我有32×32矩陣，有36個工作項目等get_global_id（0）從0 - > 35我想，和大小= MATRIX_DIM/4 = 8。
__kernel void transpose(__global float4 *g_mat,
__local float4 *l_mat, uint size) {
__global float4 *src, *dst;
/* Determine row and column location */
int col = get_global_id(0);
int row = 0;
while(col >= size) {
col -= size--;
row++;
}
col += row;
size += row;
/* Read source block into local memory */
src = g_mat + row * size * 4 + col;
l_mat += get_local_id(0)*8;
在clEnqueueNDRangeKernel調用，ARG local_work_size設置爲NULL根據手動裝置，它讓編譯器什麼的數字出來：
local_work_size can also be a NULL value in which case the OpenCL implementation will determine how to be break the global work-items into appropriate work-group instances.
但我不明白的乘法由8，這給了一個地址偏移到我認爲工作組的本地內存。有人可以解釋這個嗎？
l_mat[0] = src[0];
l_mat[1] = src[size];
l_mat[2] = src[2*size];
l_mat[3] = src[3*size];
/* Process block on diagonal */
if(row == col) {
src[0] =
(float4)(l_mat[0].x, l_mat[1].x, l_mat[2].x, l_mat[3].x);
src[size] =
(float4)(l_mat[0].y, l_mat[1].y, l_mat[2].y, l_mat[3].y);
src[2*size] =
(float4)(l_mat[0].z, l_mat[1].z, l_mat[2].z, l_mat[3].z);
src[3*size] =
(float4)(l_mat[0].w, l_mat[1].w, l_mat[2].w, l_mat[3].w);
}
/* Process block off diagonal */
else {
/* Read destination block into local memory */
dst = g_mat + col * size * 4 + row;
l_mat[4] = dst[0];
l_mat[5] = dst[size];
l_mat[6] = dst[2*size];
l_mat[7] = dst[3*size];
/* Set elements of destination block */
dst[0] =
(float4)(l_mat[0].x, l_mat[1].x, l_mat[2].x, l_mat[3].x);
dst[size] =
(float4)(l_mat[0].y, l_mat[1].y, l_mat[2].y, l_mat[3].y);
dst[2*size] =
(float4)(l_mat[0].z, l_mat[1].z, l_mat[2].z, l_mat[3].z);
dst[3*size] =
(float4)(l_mat[0].w, l_mat[1].w, l_mat[2].w, l_mat[3].w);
/* Set elements of source block */
src[0] =
(float4)(l_mat[4].x, l_mat[5].x, l_mat[6].x, l_mat[7].x);
src[size] =
(float4)(l_mat[4].y, l_mat[5].y, l_mat[6].y, l_mat[7].y);
src[2*size] =
(float4)(l_mat[4].z, l_mat[5].z, l_mat[6].z, l_mat[7].z);
src[3*size] =
(float4)(l_mat[4].w, l_mat[5].w, l_mat[6].w, l_mat[7].w);
}
}
「每個線程都需要8個float4s」你是指在讀入本地內存時它是如何運行l_mat [0] - > l_mat [7]的權利？這是有道理的......但我認爲每個'float4'都是16個字節，而我們的內存是字節可尋址的 - 稍等一下，l_mat被聲明爲一個指向float4s的指針，所以我猜想對齊是被照顧的。那麼，如果這就是我所知道的，但這是棘手的東西！ – JDS
@YoungMoney：這裏沒有對齊問題，這只是通常的C指針算術，它以字節爲單位對底層元素的大小進行了完全的抽象。 – Pragmateek