CUDA केवल एक चर के लिए L1 कैश अक्षम करें

Question

क्या CUDA 2.0 डिवाइस पर केवल एक विशिष्ट चर के लिए L1 कैश अक्षम करने का कोई तरीका है? मुझे पता है कि कोई भी सभी मेमोरी ऑपरेशंस के लिए ध्वज -Xptxas -dlcm=cgnvcc को संकलित समय में संकलित समय पर L1 कैश अक्षम कर सकता है। हालांकि, मैं केवल विशिष्ट ग्लोबल वैरिएबल पर मेमोरी पढ़ने के लिए कैश को अक्षम करना चाहता हूं ताकि शेष मेमोरी एल 1 कैश के माध्यम से जा सके।

मैंने खोज में वेब पर किया है, एक संभावित समाधान पीटीएक्स असेंबली कोड के माध्यम से है।

Answer 1

जैसा कि बताया जा ऊपर आप इनलाइन PTX उपयोग कर सकते हैं, यहाँ एक उदाहरण है:

__device__ __inline__ double ld_gbl_cg(const double *addr) { 
    double return_value; 
    asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); 
    return return_value; 
} 

आप आसानी से .f32 (नाव) या .s32 (पूर्णांक) के लिए .f64 स्वैप करके इस भिन्न हो सकते हैं आदि, "_ एफ "(फ्लोट) या" = आर "(int) etc के लिए return_value" = d "की बाधा। ध्यान दें कि (addr) - "l" से पहले अंतिम बाधा - 64 बिट एड्रेसिंग का संकेत देती है, यदि आप 32 बिट एड्रेसिंग का उपयोग कर रहे हैं, तो यह "आर" होना चाहिए।

Answer 2

इनलाइन पीटीएक्स का उपयोग चर को लोड और स्टोर करने के लिए किया जा सकता है। ld.cg और st.cg निर्देश केवल L2 में डेटा कैश करते हैं। कैश ऑपरेटरों को PTX ISA 2.3 दस्तावेज़ के अनुभाग 8.7.8.1 कैश ऑपरेटरों में वर्णित किया गया है। निर्देश या ब्याज एलडी और सेंट हैं। इनलाइन पीटीएक्स का वर्णन Using Inline PTX Assembly in CUDA में किया गया है।

Answer 3

यदि आप चर को volatile घोषित करते हैं, तो यह केवल फर्मि जीपीयू पर एल 2 कैश में कैश किया जाएगा। ध्यान दें कि कुछ कंपाइलर अनुकूलन, जैसे बार-बार लोड को हटाने, अस्थिर चर पर प्रदर्शन नहीं किए जाते हैं क्योंकि संकलक मानते हैं कि उन्हें किसी अन्य थ्रेड द्वारा लिखा जा सकता है।