2012-01-29 11 views
6

Tôi đang cố gắng hiểu kiến ​​trúc của các thiết bị OpenCL như GPU và tôi không thấy tại sao có giới hạn rõ ràng về số lượng mục công việc trong nhóm làm việc cục bộ, tức là hằng số CL_DEVICE_MAX_WORK_GROUP_SIZE. Dường như với tôi rằng điều này cần được thực hiện bởi trình biên dịch, tức là nếu một hạt nhân (một chiều đơn giản) được thực thi với kích thước nhóm làm việc cục bộ 500 trong khi mức tối đa vật lý của nó là 100, và hạt nhân tìm kiếm ví dụ như thế này:Tại sao có CL_DEVICE_MAX_WORK_GROUP_SIZE?

__kernel void test(float* input) { 
    i = get_global_id(0); 
    someCode(i); 
    barrier(); 
    moreCode(i); 
    barrier(); 
    finalCode(i); 
} 

sau đó nó có thể được chuyển đổi tự động để thực hiện một công việc với quy mô nhóm 100 trên kernel này:

__kernel void test(float* input) { 
    i = get_global_id(0); 
    someCode(5*i); 
    someCode(5*i+1); 
    someCode(5*i+2); 
    someCode(5*i+3); 
    someCode(5*i+4); 
    barrier(); 
    moreCode(5*i); 
    moreCode(5*i+1); 
    moreCode(5*i+2); 
    moreCode(5*i+3); 
    moreCode(5*i+4); 
    barrier(); 
    finalCode(5*i); 
    finalCode(5*i+1); 
    finalCode(5*i+2); 
    finalCode(5*i+3); 
    finalCode(5*i+4); 
} 

Tuy nhiên, có vẻ như điều này không được thực hiện theo mặc định. Tại sao không? Có cách nào để làm cho quá trình này tự động (khác hơn là viết một trình biên dịch trước cho nó bản thân mình)? Hoặc là có một vấn đề nội tại mà có thể làm cho phương pháp của tôi thất bại trên một số ví dụ (và bạn có thể cho tôi một)?

+0

Nó sẽ đặt năm cuộc gọi chức năng theo trình tự, phải không? Vì vậy, nó sẽ là một loại rơi trở lại, chỉ. Ngoài ra bạn sẽ phải chắc chắn để pad kích thước làm việc của bạn đến một bội số của năm. Chỉ cần tự hỏi .. – rdoubleui

Trả lời

4

Tôi nghĩ rằng nguồn gốc của CL_DEVICE_MAX_WORK_GROUP_SIZE nằm trong triển khai phần cứng cơ bản.

Nhiều luồng đang chạy đồng thời trên các đơn vị tính toán và mỗi chủ đề cần giữ trạng thái (đối với cuộc gọi, jmp, v.v.). Hầu hết các triển khai sử dụng một ngăn xếp cho điều này và nếu bạn nhìn vào gia đình AMD Evergreen, chúng là một giới hạn phần cứng cho số lượng các mục stack có sẵn (mỗi mục stack có subentries). Mà trong bản chất giới hạn số lượng các chủ đề mỗi đơn vị máy tính có thể xử lý đồng thời.

Đối với trình biên dịch có thể làm điều này để làm cho nó có thể. Nó có thể làm việc nhưng hiểu rằng nó sẽ có nghĩa là biên dịch lại hạt nhân một lần nữa. Mà không phải lúc nào cũng có thể. Tôi có thể tưởng tượng các tình huống mà các nhà phát triển đổ hạt nhân đã biên dịch cho mỗi nền tảng theo định dạng nhị phân và gửi nó với phần mềm của họ chỉ vì các lý do "không phải mã nguồn mở".

+0

+1 cho hạt nhân nhị phân, sẽ là một yếu tố hạn chế nhất định. – rdoubleui

0

Các hằng số được truy vấn từ thiết bị bởi trình biên dịch để xác định kích thước nhóm công việc phù hợp tại thời gian biên dịch (nơi biên dịch khóa học liên quan đến biên dịch hạt nhân). Tôi có thể làm cho bạn sai, nhưng có vẻ như bạn đang nghĩ đến việc tự mình thiết lập những giá trị đó, điều đó sẽ không xảy ra.

Trách nhiệm nằm trong mã của bạn để truy vấn các khả năng của hệ thống được chuẩn bị cho bất kỳ phần cứng nào nó sẽ chạy.