CUDA編程中的Struture of Arrays問題

最近在寫CUDA的程式,遇到一些關於設計的問題,想在這裡請問大家。

場景:
公司裡有一群員工(Employee)的資料,每個員工資料包含了id、薪水、績效和上司。公司年終的時候會比較員工和其上司的績效(績效除以薪水的比例)來給定評語。

直觀上我們可以用pointer-based和array of structure (AOS)的實作

// 方法一 (pointer-based + AOS)
struct Employee{
   size_type id;
   value_type salary;
   value_type performance;
   Employee* supervisor;
};

const char* const judge(Employee* emp){
  cp_value_emp = emp->performance / emp->salary;

  Employee* sup = emp->superviosr;
  cp_value_sup = sup->performance / sup->salary;

  return cp_value_emp > cp_value_sup ? "super!" : "bad!";
};

//array of structure (AOS)
Employee* employees;

然而在GPU計算中,考慮需要先把資料從CPU memory複製到GPU memory,這時方法一可能會不太方便,其中一個可能的解決方法是用index-based和structure of array (SOA)的實作

//方法二 (index-based + SOA)
struct Employees{
  size_type* ids;
  value_type* salaries;
  value_type* performances;
  size_type* supervisor_indices;
};

const char* const judge(Employees emps, size_type index){
  cp_value_emp = emps.performances[index] / emps.salaries[index];

  size_type sup_index = emps.superviosrs[index];
  cp_value_sup = emps.performances[sup_index] / emps.salaries[sup_index];

  return (cp_value_emp > cp_value_sup) ? "super!" : "bad!";
};

// structure of array (SOA)
Employees employees;

基於方法二,我們也可以不封裝成structure,將各個array分開獨立來看

//方法三 (index-based + separated arrays)
size_type* ids;
value_type* salaries;
value_type* performances;
size_type* supervisor_indices;

const char* const judge(value_type* salaries, value_type* performances, size_type* supervisor_indices, size_type index){
  cp_value_emp = performances[index] / salaries[index];

  size_type sup_index = superviosrs[index];
  cp_value_sup = performances[sup_index] / salaries[sup_index];

  return (cp_value_emp > cp_value_sup) ? "super!" : "bad!";
};

不知道在CUDA實作中,方法二和方法三各有什麼優缺點呢?
目前的想法是,方法二可能會需要在CUDA額外宣告Employees這個structure,不確定會不會造成額外的性能開銷,方法三的缺點是函數的argument會變的很冗長,降低可讀性。


根據 @yhmtsai 的回覆,延伸出的方法四

// 方法四 替使用者和開發者提供不同的函數介面

struct Employees{
  size_type num_employee;
  size_type* ids;
  value_type* salaries;
  value_type* performances;
  size_type* supervisor_indices;
};

// 供使用者使用的函數介面
const char* const judge(Employees emps, size_type index){
  size_type nemp = emps.num_employee;
  value_type* d_salaries;
  value_type* d_performances;
  value_type* d_supervisor_indices;
  cudaMalloc(&d_salaries, nemp);
  cudaMalloc(&d_performances, nemp);
  cudaMalloc(&d_supervisor_indices, nemp);
  cudaMemcpy(d_salaries, emps.salaries, cudaMemcpyHostToDevice);
  cudaMemcpy(d_performances, emps.performances, nemp, cudaMemcpyHostToDevice);
  cudaMemcpy(d_supervisor_indices, emps.supervisors, nemp, cudaMemcpyHostToDevice);
  return __gpu_judge(d_salaries, d_performances, d_supervisor_indices, index);
}

// 開發者才會接觸到的內部GPU函數
__global__ const char* const __gpu_judge(value_type* d_salaries, value_type* d_performances, size_type* d_supervisor_indices, size_type index){
  ...
}

argument 冗長個人覺得相對還好,只要變數名稱有弄好的話
方法三應該會相對簡單很多
另外對於 __restrict__ 關鍵字的使用,也會比較方便。
如果放到 structure 裏頭,要讓 compiler 要使用到 __restrict__ 需要下額外功夫
方法二應該是直接傳入 Employees 就好,而非 Employees *?

AoS 的問題是蠻容易造成 uncoalesced memory read/write (非連續記憶體) 導致效能低落。
例如當要讀取 emp[threadIdx.x]->id 時,因為記憶體的擺放,各個的 id 中間還會有 salary, performance, superviosr,所以 CUDA 要多花幾次記憶體讀取才能把 id 給準備好。

備註: 我印象中也有 SoAoS 或者 AoSoA 的設計,但這個就更複雜,但基本上都是使用記憶體越少刺越好

1個讚

感謝回答,原來還有__restrict__這個關鍵字可以用。
我目前也是覺得方法二有點麻煩,會需要一些額外的功夫,而想採用方法三。
根據你的回覆我又延伸出了方法四,替使用者端(CPU)和開發者端(GPU)提供不同的介面。這個方法如何呢?

這也是我們常用的方式
除非是像 CUDA 那種直接提供 GPU 的介面
通常是用 class 或者 structure 包好給使用者使用,當需要用到 GPU 時,再寫怎麼將成員變數傳進 GPU 的介面(這邊就不會給使用者使用到)

1個讚