OpenCL的標VS矢量

我有簡單的內核：OpenCL的標VS矢量

__kernel vecadd(__global const float *A, 
       __global const float *B, 
       __global float *C) 
{ 
    int idx = get_global_id(0); 
    C[idx] = A[idx] + B[idx]; 
}

爲什麼當我改變浮到個float4，內核運行慢30％以上？

所有教程說，使用矢量類型的加速計算...

在主機端，內存爲alocated參數float4變量爲16個字節對齊，並global_work_size爲clEnqueueNDRangeKernel小4倍。

內核在AMD HD5770 GPU，AMD-APP-SDK-v2.6上運行。

爲CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT設備信息返回4

編輯：
global_work_size = 1024 * 1024（以及更大）
local_work_size = 256
時間使用CL_PROFILING_COMMAND_START和CL_PROFILING_COMMAND_END測量。

對於較小的global_work_size（浮點型爲8196/float4爲2048），向量化版本更快，但我想知道爲什麼？

來源

2012-01-19 ldanko

全球工作規模和工作組規模的價值是什麼？你什麼時候測量，以及如何？ –

全局工作大小= 1024 * 1024本地工作大小= 256，我使用CL_PROFILING_COMMAND_START和CL_PROFILING_COMMAND_END測量clEnquueNDRangeKernel的時間。對於較小的global_work_size（8196代表float/2048代表float4），向量化版本更快，但我想知道，爲什麼？ – ldanko

較小和較大的工作量之間的差異可能是由於您的緩存不斷。所以2個問題： 1）如果你刪除了const，它對於小而言還是比較快，對於較大的則是較慢？ 2）如果你去中間的某個地方，比如說浮動的65536和浮動的16384，那麼會發生什麼？ – user1111929

我不知道你提到的教程是什麼，但它們必須是舊的。 ATI和NVIDIA現在都使用標量GPU架構至少五年。現在，在您的代碼中使用矢量僅用於語法方便，與普通標量代碼相比，它沒有性能優勢。事實證明，對於GPU，標量架構比矢量更好 - 它更好地利用硬件資源。

來源

2012-04-19 11:03:00 lucho

我不確定爲什麼矢量對於你來說要慢得多，而不瞭解工作組和全局大小。我希望它至少有相同的表現。

如果它適合你的內核，你可以從C開始使用A中的值嗎？這將減少33％的內存訪問。也許這適用於你的情況？

__kernel vecadd(__global const float4 *B, 
       __global float4 *C) 
{ 
    int idx = get_global_id(0); 
    C[idx] += B[idx]; 
}

另外，你是否厭倦了讀取私人載體的值，然後添加？或者可能是兩種策略。

__kernel vecadd(__global const float4 *A, 
       __global const float4 *B, 
       __global float4 *C) 
{ 
    int idx = get_global_id(0); 
    float4 tmp = A[idx] + B[idx]; 
    C[idx] = tmp; 
}

來源

2012-01-20 16:00:26 mfa

OpenCL的標VS矢量

回答

相關問題