การลด OpenCL จากส่วนตัวไปสู่ระดับท้องถิ่นและระดับโลก?

เคอร์เนลต่อไปนี้คำนวณสนามความดันเสียง โดยแต่ละเธรดจะคำนวณเป็นอินสแตนซ์ส่วนตัวของเวกเตอร์ pressure ซึ่งจะต้องถูกสรุปลงในหน่วยความจำส่วนกลาง ฉันค่อนข้างแน่ใจว่าโค้ดที่คำนวณ pressurevector นั้นถูกต้อง แต่ฉันยังคงประสบปัญหาในการทำให้ได้ผลลัพธ์ตามที่คาดหวัง

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


person Emanuel Ey    schedule 19.02.2013    source แหล่งที่มา


คำตอบ (1)


ฉันสังเกตเห็นสองประเด็นที่นี่ ประเด็นใหญ่และปัญหาเล็กอีกประเด็นหนึ่ง:

  • รหัสนี้แสดงให้เห็นว่าคุณมีความเข้าใจผิดเกี่ยวกับสิ่งที่กั้นทำ สิ่งกีดขวางไม่เคยซิงโครไนซ์กับกลุ่มงานหลายกลุ่ม มันซิงโครไนซ์ภายในเวิร์กกรุ๊ปเท่านั้น CLK_GLOBAL_MEM_FENCE ทำให้ดูเหมือนว่าเป็นการซิงโครไนซ์ทั่วโลก แต่จริงๆ แล้วไม่ใช่ แฟล็กนั้นกั้นการเข้าถึงหน่วยความจำส่วนกลางของรายการงานปัจจุบันทั้งหมด การเขียนที่โดดเด่นจะสามารถสังเกตได้ทั่วโลกหลังจากมีสิ่งกีดขวางด้วยธงนี้ แต่จะไม่เปลี่ยนพฤติกรรมการซิงโครไนซ์ของอุปสรรค ซึ่งอยู่ในขอบเขตของกลุ่มงานเท่านั้น ไม่มีการซิงโครไนซ์ทั่วโลกใน OpenCL นอกเหนือจากการเปิดตัว NDRange หรืองานอื่น
  • การวนซ้ำครั้งแรกจะทำให้รายการงานหลายรายการเขียนทับการคำนวณของกันและกัน การสร้างดัชนีของ pressure_local ด้วย iHyd จะดำเนินการโดยแต่ละรายการงานที่มี iHyd เดียวกัน สิ่งนี้จะสร้างผลลัพธ์ที่ไม่ได้กำหนด

หวังว่านี่จะช่วยได้

person boiler96    schedule 19.02.2013
comment
ขอบคุณสำหรับคำตอบ และขออภัยสำหรับการตอบกลับล่าช้า: จริง ๆ แล้วมีอยู่ช่วงหนึ่งที่ฉันได้วางลูปสุดท้ายไว้ในเคอร์เนลที่แยกจากกัน ฉันปล่อยให้โค้ดนั่งสักพักแล้วลืมไปว่าทำไมฉันถึงทำแบบนั้น - ขอบคุณที่เตือนฉัน :p ฉันจะแยกมันออกเป็นเคอร์เนลแยกอีกครั้งแล้วดูลูปแรกอีกครั้ง - person Emanuel Ey; 24.02.2013
comment
เฮ้ ฉันประสบปัญหาเดียวกัน และสงสัยว่าคุณจัดการมันได้อย่างไร - person Eric; 13.03.2015