मैं पंजीकरण उपयोग को ट्रैक करने की कोशिश कर रहा हूं और एक दिलचस्प परिदृश्य में आया हूं। निम्नलिखित स्रोत पर विचार करें:cuda कर्नेल रजिस्टर उपयोग को ट्रैक करना
#define OL 20
#define NHS 10
__global__ void loop_test(float ** out, const float ** in,int3 gdims,int stride){
const int idx = blockIdx.x*blockDim.x + threadIdx.x;
const int idy = blockIdx.y*blockDim.y + threadIdx.y;
const int idz = blockIdx.z*blockDim.z + threadIdx.z;
const int index = stride*gdims.y*idz + idy*stride + idx;
int i = 0,j =0;
float sum =0.f;
float tmp;
float lf;
float u2, tW;
u2 = 1.0;
tW = 2.0;
float herm[NHS];
for(j=0; j < OL; ++j){
for(i = 0; i < NHS; ++i){
herm[i] += in[j][index];
}
}
for(j=0; j<OL; ++j){
for(i=0;i<NHS; ++i){
tmp = sum + herm[i]*in[j][index];
sum = tmp;
}
out[j][index] = sum;
sum =0.f;
}
}
स्रोत पर एक पक्ष नोट के रूप में - बस कहते हैं - चल रहा है योग मैं कर सकता + =, लेकिन कैसे बदल रहा है कि प्रभाव के उपयोग रजिस्टर (लगता है यह नहीं है के साथ खेल रहा था एक अतिरिक्त mov निर्देश)। इसके अतिरिक्त यह स्रोत 3 डी स्पेस में मैप किए गए मेमोरी तक पहुंचने के लिए उन्मुख है।
रजिस्टरों की गणना करना ऐसा लगता है कि 22 रजिस्ट्रार हैं (मुझे लगता है कि एक फ्लोट [एन] एन + 1 रजिस्ट्रार लेता है - अगर मैं गलत हूं तो कृपया मुझे सही करें) घोषणाओं के आधार पर।
nvcc -cubin -arch=sm_20 -Xptxas="-v" src/looptest.cu
पैदावार:
हालांकि साथ संकलन
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 25 registers, 72 bytes cmem[0]
ठीक है तो नंबर क्या 'उम्मीद' है कि अलग है। साथ संकलित इसके अतिरिक्त यदि:
nvcc -cubin -arch=sm_13 -Xptxas="-v" src/looptest.cu
रजिस्टर उपयोग है दूर कम - 8 (चल बिन्दु गणित मानकों आईईईई के लिए जाहिरा तौर पर sm_13 से sm_20 में मजबूत पालन के कारण?) सटीक होना करने के लिए:
ptxas info : Compiling entry function '_Z9loop_testPPfPPKfS2_4int3i' for 'sm_13'
ptxas info : Used 17 registers, 40+16 bytes smem, 8 bytes cmem[1]
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 28 registers, 72 bytes cmem[0]
मुझे पता है कि जहां रजिस्टर किया जा रहा है करना चाहते हैं:
एक अंतिम नोट के रूप में, 40 के लिए मैक्रो राजभाषा बदलने के लिए, और अचानक खाया, और मैंने जो जोड़े अवलोकन किए हैं, उसके परिणाम क्या हैं।
मेरे पास एक cuobjdump के माध्यम से असेंबली के साथ पर्याप्त अनुभव नहीं है - उत्तर निश्चित रूप से वहां दफनाया जाता है - शायद कोई मुझे इस बारे में बता सकता है कि मुझे क्या देखना चाहिए या मुझे कैसे संपर्क करना है इस बारे में एक गाइड दिखाएं असेंबली डंप
यह हो सकता है कि आपके छोरों मूल्य 20 से राजभाषा के लिए संकलक द्वारा unrolled किया गया और 40 के लिए उतारना नहीं किया? –
मुझे लगता है कि अश्विन की टिप्पणी सही है। सीयूडीए सी प्रोग्रामिंग गाइड में उल्लिखित अनुसार, आपको वार्प लेवल एडिशन मामलों के माध्यम से अपने लूपिंग रम्स को फ़्लैट करने पर विचार करना चाहिए। http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Programming_Guide.pdf –
मैं बहुत विश्वास है कि रजिस्टर गिनती में अंतर चल बिन्दु, या पाश के साथ कुछ नहीं करना होगा अनियंत्रित या अब तक किसी और का उल्लेख किया गया है। याद रखें कि sm_20 आंतरिक रूप से 64 बिट आर्किटेक्चर है और sm_13 32 बिट आर्किटेक्चर है। इसका मतलब है कि पॉइंटर्स को sm_12 की तुलना में sm_20 के लिए संकलित रजिस्टर पदचिह्न से दोगुना है। – talonmies