Pengurangan OpenCL dari swasta ke lokal lalu global?

Kernel berikut menghitung bidang tekanan akustik, dengan masing-masing thread menghitung instance pribadinya dari vektor pressure, yang kemudian perlu dijumlahkan ke dalam memori global. Saya cukup yakin kode yang menghitung pressurevector sudah benar, tetapi saya masih kesulitan membuat ini memberikan hasil yang diharapkan.

int gid       = get_global_id(0);
int lid       = get_local_id(0);
int nGroups   = get_num_groups(0);
int groupSize = get_local_size(0);
int groupID   = get_group_id(0);

/* Each workitem gets private storage for the pressure field.
 * The private instances are then summed into local storage at the end.*/
private float2    pressure[HYD_DIM_TOTAL];
local   float2    pressure_local[HYD_DIM_TOTAL];

/* Code which computes value of 'pressure' */

//wait for all workgroups to finish accessing any memory
barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);

/// sum all results in a workgroup into local buffer:
for(i=0; i<groupSize; i++){

    //each thread sums its own private instance into the local buffer
    if (i == lid){
        for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){
            pressure_local[iHyd] += pressure[iHyd];
        }
    }
    //make sure all threads in workgroup get updated values of the local buffer
    barrier(CLK_LOCAL_MEM_FENCE);
}

/// copy all the results into global storage
//1st thread in each workgroup writes the group's local buffer to global memory
if(lid == 0){
    for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){
        pressure_global[groupID +nGroups*iHyd] = pressure_local[iHyd];
    }
}

barrier(CLK_GLOBAL_MEM_FENCE);

/// sum the various instances in global memory into a single one
// 1st thread sums global instances
if(gid == 0){

    for(iGroup=1; iGroup<nGroups; iGroup++){

        //we only need to sum the results from the 1st group onward
        for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){

            pressure_global[iHyd] += pressure_global[iGroup*HYD_DIM_TOTAL +iHyd];
            barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }
}

Beberapa catatan mengenai dimensi data: Jumlah total thread akan bervariasi antara 100 dan 2000, namun terkadang berada di luar interval ini.
groupSizeakan bergantung pada perangkat keras, namun saat ini saya menggunakan nilai antara 1(cpu) dan 32(gpu ).
HYD_DIM_TOTAL diketahui pada waktu kompilasi dan bervariasi antara 4 dan 32 (umumnya, namun tidak harus, berupa pangkat 2).

Apakah ada yang salah dengan kode pengurangan ini?

PS: Saya menjalankan ini pada i7 3930k dengan AMD APP SDK 2.8 dan pada NVIDIA GTX580.


amd-processor gpu opencl reduction gpgpu
person Emanuel Ey    schedule 19.02.2013    source sumber

Jawaban (1)


arrow_upward
4
arrow_downward

Saya melihat dua masalah di sini, satu besar, satu lagi lebih kecil:

  • Kode ini menunjukkan bahwa Anda memiliki kesalahpahaman tentang fungsi penghalang. Sebuah penghalang tidak pernah disinkronkan di beberapa kelompok kerja. Itu hanya melakukan sinkronisasi dalam kelompok kerja. CLK_GLOBAL_MEM_FENCE membuatnya tampak seperti sinkronisasi global, padahal sebenarnya tidak. Bendera itu hanya memagari semua akses item kerja saat ini ke memori global. Jadi tulisan yang luar biasa akan dapat diamati secara global setelah adanya penghalang dengan bendera ini. Tapi itu tidak mengubah perilaku sinkronisasi penghalang, yang hanya pada lingkup kelompok kerja. Tidak ada sinkronisasi global di OpenCL, selain meluncurkan NDRange atau Tugas lain.
  • Perulangan for pertama menyebabkan beberapa item kerja menimpa perhitungan satu sama lain. Pengindeksan pressure_local dengan iHyd akan dilakukan oleh setiap item pekerjaan dengan iHyd yang sama. Ini akan menghasilkan hasil yang tidak terdefinisi.

Semoga ini membantu.

person boiler96    schedule 19.02.2013
comment
Terima kasih atas jawabannya, dan maaf atas balasan saya yang terlambat: Saya sebenarnya pernah menempatkan loop terakhir di kernel terpisah. Saya membiarkan kodenya diam sebentar dan lupa mengapa saya melakukan itu -terima kasih telah mengingatkan saya :p Saya akan membaginya menjadi kernel terpisah lagi dan melihat lagi loop pertama - person Emanuel Ey; 24.02.2013
comment
Hai, saya mengalami masalah yang sama dan bertanya-tanya bagaimana Anda berhasil mengatasinya. - person Eric; 13.03.2015