क्या यह ओपनसीएल कोड अनुकूलित किया जा सकता है?

Question

मैं एक विशेष मैट्रिक्स समारोह के लिए OpenCL कोड के एक टुकड़े पर काम कर रहा हूँ: एक Dx1 वेक्टर v, दो DxD मैट्रिक्स A और B और एक निरंतर c के लिए, 1xD वेक्टर r लौट जहां r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

नीचे क्या मैं इतना है दूर, लेकिन यह अजीब धीमी गति से चलाता है। संक्षेप में एक संस्करण जो DxD मैट्रिक्स देता है लगभग दस गुना तेज है। इसे PyOpenCL से बुलाया जाता है यदि इससे कोई फर्क पड़ता है।

क्या कुछ गलत है? क्या इसे अनुकूलित किया जा सकता है?

#define D 1000 
... 

    __kernel void element_mult(
     __global float *result, 
     __global const float *vector, 
     __global const float *matrix, 
     __global const float *matrix2, 
     const float factor) 
     { 
     int y = get_global_id(1); 
     float sum = 0; 
     for(int k = 0; k < D; k++) 
     { 
      sum += vector[k] * matrix[(y*D) + k] 
      * matrix2[(y*D) + k ]; 
     } 
     result[y] = sum * factor; 
     } 

चीयर्स!

Answer 1

अनुकूलन # 1: वेक्टर __local बनाओ।

इस पर मेरा पहला पास प्रदर्शन में एक सभ्य सुधार मिला। मैंने देखा कि प्रत्येक वेक्टर [के] कुल डी बार पढ़ा जाता है, इसलिए मैंने इसे __local में कॉपी किया। यह केवल इसलिए संभव है क्योंकि डी इसे अनुमति देने के लिए काफी छोटा है। जैसा कि आपके ऊपर है, कर्नेल एक भयानक एएलयू से पीड़ित है: 5870 और 6970 जीपीएस दोनों पर 0.08 का fetch अनुपात। यहां तक ​​कि धीमे gpus अभी भी स्मृति पहुंच पर इंतजार कर रहे हैं।

#define D 1000 
    __kernel void element_mult(
    __global float *result, 
    __global const float *vector, 
    __global const float *matrix, 
    __global const float *matrix2, 
    const float factor) 
    { 
     int y = get_global_id(0); 
     float sum = 0; 

     __local float vectCopy[D]; 
     int ls = get_local_size(0); 
     int lid = get_local_id(0); 
     for(int i=0;i<D;i+=ls){ 
      vectCopy[i+lid] = vector[i+lid]; 
     } 
     mem_fence(CLK_LOCAL_MEM_FENCE); 

     for(int k = 0; k < D; k++) 
     { 
      sum += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ]; 
     } 
     result[y] = sum * factor; 
    } 

इस बदलाव के साथ

, एपीपी प्रोफाइलर एक नया ALU दिखा रहा है: 5870 और 6970 GPUs के लिए 0.20 के अनुपात से लाएं। 1513 -> 1034, और 1261 -> 861 से औसत कार्ड एक ही कार्ड पर बदल गए। कम अंत gpus अब लाने के बजाय एएलयू द्वारा बंधे हैं। (4: 1 अनुपात से अधिक)

ओपिमनाइज़ेशन # 2: प्रत्येक परिणाम की गणना करें [y] एक संपूर्ण कार्य समूह का उपयोग करके।

आपको यह आईडी करना होगा डी डी बहुत बड़ा (100k +) था। विचार एक समय में परिणाम के एक तत्व की गणना करने के लिए कार्य समूह का उपयोग करके सर्वोत्तम स्मृति पहुंच पैटर्न प्राप्त करना है। मैंने एलएस (स्थानीय आकार) को यहां 64 होने के लिए परिभाषित किया है, क्योंकि यह मेरे हार्डवेयर, साथ ही साथ अधिकांश विक्रेताओं पर भी काम करता है। आपके द्वारा होस्ट-साइड से उपयोग किए जाने वाले कार्यसमूह आकार को तब तक 64 होना होगा जब तक आप उस परिभाषा को बदल नहीं देते। इसे __local के रूप में योग [एलएस] भंडारण बनाने के लिए परिभाषित करने की आवश्यकता है, और मुझे अपने कर्नेल में परिवर्तनीय आकार के __local vars को गुजरना पसंद नहीं है।

परिणाम: 5870 एएलयू: fetch = 0.59: 1, औसत = 708। 6 9 70 एएलयू: fetch = 0.72, औसत = 5 9 0। एपीपी प्रोफाइलर के अनुसार, यह आपकी मूल सूची के जितना तेज़ है।

#define D 1000 
#define ls 64 
__kernel void element_mult(
__global float *result, 
__global const float *vector, 
__global const float *matrix, 
__global const float *matrix2, 
const float factor) 
{ 
    __local float vectCopy[D]; 
    int lid = get_local_id(0); 
    for(int i=0;i<D;i+=ls){ 
     vectCopy[i+lid] = vector[i+lid]; 
    } 
    mem_fence(CLK_LOCAL_MEM_FENCE); 

    int ng = get_num_groups(0); 
    int gid = get_group_id(0); 
    int y, k; 
    __local float sum[ls]; 
    for(y = gid; y < D; y+=ng){ 
     for(k = lid; k < D; k+=ls) 
     { 
      sum[lid] += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ]; 
     } 
     if(lid==0){ 
      result[y] = sum[0]; 
      for(k=1;k<ls;k++){ 
       result[y] += sum[k]; 
      } 
      result[y] *= factor; 
     } 
     mem_fence(CLK_LOCAL_MEM_FENCE); 
    } 
} 

संपादित करें: एपीपी प्रोफाइलर = एएमडी एपीपी KernelAnalyzer