2013-10-09 79 views
1

我正在讀這篇文章http://eprints.dcs.warwick.ac.uk/1694/1/miniMD_opencl.pdf關於使用OpenCl的迷你分子動力學應用。代碼位於此處。
我遇到了內核如何實現。我不明白的是這OpenCl中的標量和矢量內核

#if defined(SCALAR_KERNELS) 
__kernel void f_clear(
    __global float* f, 
    __const int nall) { 

    for (unsigned i = get_global_id(0)+1; i <= nall; i += get_global_size(0)) { 
     const int i4 = i << 2; 
     f[i4+0] = 0.0f; 
     f[i4+1] = 0.0f; 
     f[i4+2] = 0.0f; 
     f[i4+3] = 0.0f; 
    } 

} 
#elif defined(VECTOR_KERNELS) 
__kernel __attribute__((vec_type_hint(float4))) 
void f_clear(
    __global float4* f, 
    __const int nall) { 

    const float4 zeroes = (float4) (0.0f, 0.0f, 0.0f, 0.0f); 
    for (unsigned i = get_global_id(0)+1; i <= nall; i += get_global_size(0)) { 
     f[i] = zeroes; 
    } 

} 
#endif 

是假設VECTOR_KERNELSSCALAR_KERNELS對應GPU和MIC設備,但不能肯定。
這是否與MIMD SIMD指令或多核和矢量編程有關?
現在還有使用矢量類型的真正優勢嗎?
最後,我真的不知道兩個for循環做什麼和他們的目的。
爲什麼不只是做f[get_global_id(0)]
謝謝,
Éric。

回答

1

一些像AMD,ATI和Intel這樣的設備對支持矢量類型非常好。這些向量是SIMD,如果可能的話,使用它更快。 NVIDIA在支持OpenCL的載體方面似乎不太好(至少我測試過的那些)。

這兩個循環似乎都會清除大小爲nall的全局內存塊。

3

標量和矢量在OpenCL中執行相同的事情只是不同的方法。 但是向量是要走的路,因爲它們應該被編譯器(CPU或GPU/FPGA)更好地優化。這樣編譯器可以自然地將SIMD單元外接。所以,如果可以的話,請使用它們。

正如Austin所說,兩個循環都清除了全局內存大小nall

但是看代碼是非常不夠的。同一工作組中的工作項目正在訪問完全不同的全局內存區域,從而破壞了整合。這將是剛剛好很多(如你所說):

__kernel __attribute__((vec_type_hint(float4))) 
void f_clear(
    __global float4* f) { 
    f[get_global_id(0)] = (float4) (0.0f, 0.0f, 0.0f, 0.0f); 
} 

和啓動這個內核與適當的整體大小(global_size = nall),讓編譯器決定在本地工作組規模。 PS:如果我必須這樣做,我更喜歡調用clEnqueueWriteBuffer並從CPU清除內存。因爲它可以與其他內核執行並行執行。