เคอร์เนลต่อไปนี้คำนวณสนามความดันเสียง โดยแต่ละเธรดจะคำนวณเป็นอินสแตนซ์ส่วนตัวของเวกเตอร์ pressure
ซึ่งจะต้องถูกสรุปลงในหน่วยความจำส่วนกลาง ฉันค่อนข้างแน่ใจว่าโค้ดที่คำนวณ pressure
vector นั้นถูกต้อง แต่ฉันยังคงประสบปัญหาในการทำให้ได้ผลลัพธ์ตามที่คาดหวัง
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);
}
}
}
หมายเหตุบางประการเกี่ยวกับมิติข้อมูล: จำนวนเธรดทั้งหมดจะแตกต่างกันระหว่าง 100 ถึง 2,000 แต่ในบางครั้งอาจอยู่นอกช่วงเวลานี้groupSize
จะขึ้นอยู่กับฮาร์ดแวร์ แต่ขณะนี้ฉันกำลังใช้ค่าระหว่าง 1(cpu) ถึง 32(gpu) ).HYD_DIM_TOTAL
เป็นที่รู้จัก ณ เวลาคอมไพล์ และแปรผันระหว่าง 4 ถึง 32 (โดยทั่วไป แต่ไม่จำเป็นต้องเป็นกำลังของ 2)
มีอะไรผิดปกติอย่างโจ่งแจ้งกับรหัสส่วนลดนี้หรือไม่?
PS: ฉันรันสิ่งนี้บน i7 3930k พร้อม AMD APP SDK 2.8 และบน NVIDIA GTX580