OpenCL（aparapi）Radeonでの緩やかな緩和

Question

OpenCLの大きな倍精度配列の単純な削減（この場合は合計）をコード化しようとしています。私は、オンラインチュートリアルを見て、これは基本的に私の問題を解決する方法であることを発見した：

#pragma OPENCL EXTENSION cl_khr_fp64 : enable 

typedef struct This_s{ 
    __global double *nums; 
    int nums__javaArrayLength; 
    __local double *buffer; 
    __global double *res; 
    int passid; 
}This; 
int get_pass_id(This *this){ 
    return this->passid; 
} 
__kernel void run(
    __global double *nums, 
    int nums__javaArrayLength, 
    __local double *buffer, 
    __global double *res, 
    int passid 
){ 
    This thisStruct; 
    This* this=&thisStruct; 
    this->nums = nums; 
    this->nums__javaArrayLength = nums__javaArrayLength; 
    this->buffer = buffer; 
    this->res = res; 
    this->passid = passid; 
    { 
     int tid = get_local_id(0); 
     int i = (get_group_id(0) * get_local_size(0)) + get_local_id(0); 
     int gridSize = get_local_size(0) * get_num_groups(0); 
     int n = this->nums__javaArrayLength; 
     double cur = 0.0; 
     for (; i<n; i = i + gridSize){ 
     cur = cur + this->nums[i]; 
     } 
     this->buffer[tid] = cur; 
     barrier(CLK_LOCAL_MEM_FENCE); 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<32){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 32)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<16){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 16)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<8){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 8)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<4){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 4)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<2){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 2)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid<1){ 
     this->buffer[tid] = this->buffer[tid] + this->buffer[(tid + 1)]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
     if (tid==0){ 
     this->res[get_group_id(0)] = this->buffer[0]; 
     } 
     return; 
    } 
} 

あなたは奇妙なthis疑問を抱いている場合、それは私が使用しaparapiの（残念ながら必要な）成果物であり、 JavaをOpenCLに変換します。

私のカーネルは正しい結果をもたらし、Nvidia Hardwareを合理的に犠牲にして、Javaの連続合計より約10倍高速です。ただし、Radeon R9 280では、単純なJavaコードと同等のパフォーマンスが得られます。

私はCodeXLでカーネルをプロファイリングしました。 MemUnitBusyがわずか6％であることがわかります。なぜそんなに低いですか？

Answer 1

OpenCLはフォルトではありませんが（直接的な）、バッファの管理はです。

私は全く同じカーネルをaparapiなしで試してみましたが、パフォーマンスは良いです。私がCL_MEM_USE_HOST_PTRを使用するとすぐに悪くなります。これは、悲しいことにaparapiを使用するときの唯一のオプションです。いくつかの "ウォーミングアップ"が実行された後でも、AMDはそのオプションでホストメモリをデバイスにコピーしていないようです。

Answer 2

aparapi.comでよりアクティブなプロジェクトに移行することを検討してください。これには、バグに対するいくつかの修正が含まれています。上にリンクした古いライブラリには、多くの追加機能とパフォーマンスの強化が含まれています。それはまた約12のリリースで中心的な場所にあります。使い方が簡単です。新しいGithub repository is here。