Perplexity ನಲ್ಲಿ CuTeDSL

Perplexity ನಲ್ಲಿ, ಎಂಭೆಡ್ಡಿಂಗ್‌ಗಳಿಂದ ಹಿಡಿದು ಟ್ರಿಲಿಯನ್-ಪ್ಯಾರಾಮೀಟರ್‌ ಎಲ್‌ಎಲ್‌ಎಂ (LLM) ಗಳವರೆಗಿನ ಮಾದರಿಗಳನ್ನು ಸರ್ವ್ ಮಾಡಲು ನಾವು ಇನ್-ಹೌಸ್ ರನ್‌ಟೈಮ್-ಆಪ್ಟಿಮೈಜ್ಡ್ ಸರ್ವಿಂಗ್ ಎಂಜಿನ್ (ROSE) ಅನ್ನು ಬಳಸುತ್ತೇವೆ. Sonar, Search ಮತ್ತು Embeddings ಸೇರಿದಂತೆ ಎಲ್ಲಾ API ಗಳ ಹಿಂದೆ ಕಾರ್ಯನಿರ್ವಹಿಸುತ್ತಿರುವ ನಮ್ಮ ರನ್‌ಟೈಮ್, ಅತ್ಯಂತ ಸೂಕ್ತವಾದ ಮಾದರಿಗಳು ಮತ್ತು ಆರ್ಕಿಟೆಕ್ಚರ್‌ಗಳನ್ನು ಅನ್ವೇಷಿಸಲು ಸಂಶೋಧಕರಿಗೆ ಅಧಿಕಾರ ನೀಡಲು ಹಾಗೂ ಉತ್ತಮ ಅಭ್ಯರ್ಥಿಗಳನ್ನು ತ್ವರಿತವಾಗಿ ಪ್ರೊಡಕ್ಷನ್‌ಗೆ ನಿಯೋಜಿಸಲು ನಿರ್ಮಿಸಲಾಗಿದೆ. ಇದನ್ನು ಸಾಧಿಸಲು, ಅತ್ಯಾಧುನಿಕ ಆರ್ಕಿಟೆಕ್ಚರ್‌ಗಳಿಗಾಗಿ ಇನ್‌ಫೆರೆನ್ಸ್ ಅನ್ನು ಅಳವಡಿಸಲು, ನಿಯೋಜಿಸಲು ಮತ್ತು ಆಪ್ಟಿಮೈಜ್ ಮಾಡಲು ನಮಗೆ ವೇಗವಾದ ಗತಿಯಲ್ಲಿ ಸಾಧ್ಯವಾಗಬೇಕು. ಈ ಬ್ಲಾಗ್ ಪೋಸ್ಟ್‌ನಲ್ಲಿ, ನಮ್ಮ ಇನ್‌ಫೆರೆನ್ಸ್ ಎಂಜಿನ್‌ನಲ್ಲಿ ಸಂಯೋಜಿಸಲಾದ CuTeDSL ಕರ್ನೆಲ್‌ಗಳು, NVIDIA GPU ಗಳಲ್ಲಿ ಮಾದರಿಗಳನ್ನು ಗರಿಷ್ಠ ಕಾರ್ಯಕ್ಷಮತೆಗೆ ತರಲು ಅಗತ್ಯವಿರುವ GPU ಕರ್ನೆಲ್‌ಗಳನ್ನು ನಿರ್ಮಿಸಲು ನಮಗೆ ಹೇಗೆ ಅವಕಾಶ ನೀಡುತ್ತವೆ ಎಂಬುದನ್ನು ನಾವು ಅನ್ವೇಷಿಸುತ್ತೇವೆ.

Perplexity ನಲ್ಲಿರುವ ವಿವಿಧ ತಂಡಗಳು ನಿರ್ಮಿಸಿದ ಕಸ್ಟಮ್ ಮಾದರಿಗಳನ್ನು, AI ಇನ್‌ಫೆರೆನ್ಸ್ ತಂಡದಿಂದ ನಿರ್ಮಿಸಲಾದ ROSE ಇನ್‌ಫೆರೆನ್ಸ್ ಎಂಜಿನ್ ಅನ್ನು ಬಳಸಿಕೊಂಡು NVIDIA Hopper ಮತ್ತು Blackwell GPU ಗಳಲ್ಲಿ ಇನ್-ಹೌಸ್ ಆಗಿ ಹೋಸ್ಟ್ ಮಾಡಲಾಗುತ್ತದೆ. ಭಾಷಾ ಡಿಕೋಡಿಂಗ್ ಮತ್ತು ವರ್ಗೀಕರಣ ಎರಡಕ್ಕೂ ಕಸ್ಟಮೈಸ್ ಮಾಡಿದ Llama ಮಾದರಿಗಳನ್ನು ಉತ್ತಮವಾಗಿ ಸರ್ವ್ ಮಾಡಲು ROSE ಅನ್ನು ಆರಂಭದಲ್ಲಿ ನಿರ್ಮಿಸಲಾಯಿತು, ಇದು NVIDIA Triton ಇನ್‌ಫೆರೆನ್ಸ್ ಸರ್ವರ್‌ನೊಂದಿಗೆ ಹೊಂದಿಕೊಳ್ಳುವ ಇಂಟರ್ಫೇಸ್ ಅನ್ನು ಹೊಂದಿದೆ. ಕಾಲಾನಂತರದಲ್ಲಿ, ಇದು ಪ್ರಮುಖ LLM ಗಳನ್ನು ಹಾಗೂ Perplexity ಯ ಉತ್ಪನ್ನಗಳಿಗೆ ಶಕ್ತಿ ತುಂಬಲು ವ್ಯಾಪಕ ಶ್ರೇಣಿಯ ತಂಡಗಳಿಗಾಗಿ ರ‍್ಯಾಂಕಿಂಗ್, ವರ್ಗೀಕರಣ, ಸ್ಕೋರಿಂಗ್ ಮತ್ತು ಇತರ ಟ್ರಾನ್ಸ್‌ಫಾರ್ಮರ್-ಆಧಾರಿತ ಮಾದರಿಗಳನ್ನು ಸರ್ವ್ ಮಾಡುವ ಪೂರ್ಣ ಪ್ರಮಾಣದ ಇನ್‌ಫೆರೆನ್ಸ್ ಎಂಜಿನ್ ಆಗಿ ಬೆಳೆಯಿತು. ROSE ನ ವಿನ್ಯಾಸದ ಗುರಿಯು ಸರಳವಾಗಿರುವುದು: ಕಾಲಾನಂತರದಲ್ಲಿ, ಗರಿಷ್ಠ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ತಲುಪಲು ಪ್ರಯತ್ನವನ್ನು ಹೂಡಿಕೆ ಮಾಡಲಾಗುತ್ತದೆ, ಆದರೆ ನಮ್ಮ ಮಾದರಿಗಳಿಗೆ ನಿರಂತರವಾಗಿ ವಿಕಸನಗೊಳ್ಳುತ್ತಿರುವ ಆರ್ಕಿಟೆಕ್ಚರ್‌ಗಳು ಮತ್ತು ಬಳಕೆ ಪ್ರಕರಣಗಳಿಗೆ ಹೊಂದಿಕೊಳ್ಳುವುದು ಮುಖ್ಯ ಆದ್ಯತೆಯಾಗಿದೆ.

ಡೇಟಾ ಪ್ರಕಾರಗಳು, ಹೆಡ್ ಡೈಮೆನ್ಷನ್‌ಗಳು, ಕಂಪೈಲ್-ಟೈಮ್ ಸ್ಥಿರಾಂಕಗಳು ಮತ್ತು ಲಾಂಚ್ ಸೆಟ್ಟಿಂಗ್‌ಗಳಾದ್ಯಂತ ಇರುವ ಕೆರ್ನಲ್ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್.

ಉನ್ನತ ಮಟ್ಟದಲ್ಲಿ, ROSE ಒಂದು ಎಂಜಿನ್ ಅನ್ನು ಹೊಂದಿದೆ, ಇದು ಮಾದರಿಯನ್ನು ಗ್ರಾಹಕರು ಸಂವಹನ ಮಾಡಬಹುದಾದ ಇಂಟರ್ಫೇಸ್‌ಗೆ ಹೊಂದಿಸುತ್ತದೆ. ಈ ಎಂಜಿನ್‌ಗಳು ಬಹಳ ಮರುಬಳಕೆಗೆ ಯೋಗ್ಯವಾಗಿದ್ದು, ವಿನಂತಿ ವೇಳಾಪಟ್ಟಿ ಮತ್ತು ಬ್ಯಾಚಿಂಗ್ ಅನ್ನು ನಿರ್ವಹಿಸುತ್ತವೆ, ಹಾಗು ಮೂಲ ಮಾದರಿಗಳಿಂದ ಡಿಕೋಡ್ ಮಾಡಿದ ಟೋಕನ್‌ಗಳು, ಎಂಬೆಡಿಂಗ್‌ಗಳು, ಅಥವಾ ಅಂಕಗಳನ್ನು ಹಿಂಪಡೆಯುತ್ತವೆ. ಸಾಧನಗಳ ಆರಂಭೀಕರಣ, ಪ್ರಕ್ರಿಯೆಗಳ ನಡುವಿನ ಹಾಗೂ ನೋಡ್‌ಗಳ ನಡುವಿನ ಸಂವಹನ ಚಾನೆಲ್‌ಗಳು, ಮತ್ತು ತೂಕಗಳ ಲೋಡಿಂಗ್—ಎಲ್ಲವನ್ನೂ ಅವುಗಳೇ ನಿರ್ವಹಿಸುತ್ತವೆ. LLMಗಳಿಗಾಗಿ, ಎಂಜಿನ್ ಬ್ಯಾಚಿಂಗ್, ಚಂಕಿಂಗ್, ಸ್ಯಾಂಪ್ಲಿಂಗ್‌ಗೆ ಹೊಣೆಗಾರವಾಗಿದೆ. ಅದು KV ಸಂಗ್ರಹಣೆ ಹಂಚಿಕೆಗೆ ಕೂಡ ಹೊಣೆಗಾರವಾಗಿದ್ದು, ಪೂರ್ಣ ಹಾಗೂ ರೇಖೀಯ attention ಎರಡಕ್ಕೂ, ಜೊತೆಗೆ prefix matchingಕ್ಕೂ ಹೊಣೆಗಾರವಾಗಿದೆ. ಎಂಬೆಡಿಂಗ್ ಎಂಜಿನ್‌ಗಳು ಸರಳವಾಗಿದ್ದು, ಅವು ಕೇವಲ ಆನ್‌ಲೈನ್ ಬ್ಯಾಚಿಂಗ್ ಅನ್ನು ನಡೆಸುತ್ತವೆ. ನಿರ್ವಹಣೆಯನ್ನು ಈ ರೀತಿಯಾಗಿ ಆಪ್ಟಿಮೈಸ್ ಮಾಡಲಾಗಿದೆ: ಮಾದರಿ ಅಕ್ಸಿಲೆರೇಟರ್‌ನಲ್ಲಿ ಸಮಾಂತರವಾಗಿ ನಡೆಯುತ್ತಿರುವಾಗಲೇ ಎಂಜಿನ್ ಮುಂದಿನ ಕೆಲಸದ ಸುತ್ತನ್ನು ಸಿದ್ಧಪಡಿಸುತ್ತದೆ.

CuTeDSL, Triton, CUDA, CUTLASS ಮತ್ತು cuBLAS ಬಳಸಿ ನಿರ್ಮಿಸಲಾದ GPU ಕರ್ನೆಲ್‌ಗಳನ್ನು ಸುತ್ತುವರೆದಿರುವ, ಮಾಡೆಲಿಂಗ್‌ಗಾಗಿ ROSE ಕಸ್ಟಮ್ ಲೇಯರ್‌ಗಳ ಸಮೃದ್ಧ ಸಂಗ್ರಹವನ್ನು ಬಹಿರಂಗಪಡಿಸುತ್ತದೆ. ಮ್ಯಾಟ್ರಿಕ್ಸ್ ಗುಣಾಕಾರ ಮತ್ತು ಗಮನ ಅನುಷ್ಠಾನಗಳಿಗಾಗಿ, ನಾವು ಪ್ರಾಥಮಿಕವಾಗಿ CUTLASS ಮತ್ತು cuBLAS ನ NVIDIA ಕರ್ನೆಲ್‌ಗಳನ್ನು ಅವಲಂಬಿಸಿದ್ದೇವೆ. ಎಂಭೆಡ್ಡಿಂಗ್, ನಾರ್ಮ್, MoE ಮತ್ತು ಆಕ್ಟಿವೇಷನ್ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಒಳಗೊಂಡಿರುವ ಪೈಪ್‌ಲೈನ್‌ನ ಉಳಿದ ಭಾಗವು ROSE ನ ಕರ್ನೆಲ್ ಲೈಬ್ರರಿಯ ಭಾಗವಾಗಿದೆ. ನಾವು ಇನ್‌ಫೆರೆನ್ಸ್ ಅನ್ನು ರನ್ ಮಾಡುವ ಹಾರ್ಡ್‌ವೇರ್ ಅನ್ನು ಉತ್ತಮವಾಗಿ ಬೆಂಬಲಿಸಲು, MoE ರೂಟಿಂಗ್, ಡಿಸ್ಪ್ಯಾಚ್ ಮತ್ತು ಕಂಬೈನ್ಗಾಗಿ ROSE ತನ್ನದೇ ಆದ ಸಾಧನಗಳ ನಡುವಿನ ಸಂವಹನ ಮೂಲಭೂತ ಅಂಶಗಳೊಂದಿಗೆ ಸಜ್ಜುಗೊಂಡಿದೆ. ಇತರ ಕರ್ನೆಲ್‌ಗಳು ತುಲನಾತ್ಮಕವಾಗಿ ಸರಳವಾಗಿದ್ದರೂ, ಅವು ರನ್ ಆಗುವ ಎಲ್ಲಾ ವಿಭಿನ್ನ ಹಿಡನ್ ಡೈಮೆನ್ಶನ್‌ಗಳು ಮತ್ತು ಸಾಧನಗಳಲ್ಲಿ ಗರಿಷ್ಠ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಸಾಧಿಸಲು ಅವುಗಳನ್ನು ಸಂಪೂರ್ಣವಾಗಿ ವಿಶೇಷಗೊಳಿಸಬೇಕು. ಹೆಚ್ಚುವರಿಯಾಗಿ, ವಿವಿಧ ಮಾದರಿಗಳ ಎಲ್ಲಾ ನಾಬ್‌ಗಳು ಮತ್ತು ಟ್ವೀಕ್‌ಗಳಿಗೆ (RMS ನಾರ್ಮ್‌ಗಾಗಿ ತೂಕ ಬಯಾಸ್ ವಿರುದ್ಧ ತೂಕ ಬಯಾಸ್ ಇಲ್ಲದಿರುವುದು, ಆಕ್ಟಿವೇಷನ್ ಕಾರ್ಯಗಳು ಇತ್ಯಾದಿ) ಅವಕಾಶ ಕಲ್ಪಿಸಲು ನಾವು ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಪೆಷಲೈಸೇಶನ್ ಅನ್ನು ಅವಲಂಬಿಸುತ್ತೇವೆ. ಪರಿಣಾಮವಾಗಿ, ಅವುಗಳನ್ನು ಕಾರ್ಯಗತಗೊಳಿಸಲು DSL ಗಳನ್ನು ಬಳಸಲು ನಾವು ನಿರ್ಧರಿಸಿದ್ದೇವೆ: ನಾವು ಆರಂಭದಲ್ಲಿ Triton ಬಳಸಿ ROSE ಅನ್ನು ನಿರ್ಮಿಸಿದ್ದೇವೆ. ಆದಾಗ್ಯೂ, ಅಂದಿನಿಂದ ನಾವು GPU ಪ್ರೋಗ್ರಾಮಿಂಗ್‌ಗಾಗಿ ಪ್ರಾಥಮಿಕ ಪರಿಸರವಾಗಿ CuTeDSL ಅನ್ನು ಅಳವಡಿಸಿಕೊಂಡಿದ್ದೇವೆ.

ಗಣಿತದ ದೃಷ್ಟಿಕೋನದಿಂದ, ಇನ್‌ಫೆರೆನ್ಸ್ ಕರ್ನೆಲ್‌ಗಳು ಸಾಮಾನ್ಯವಾಗಿ ಅತ್ಯಾಧುನಿಕವಾಗಿರುವುದಿಲ್ಲ. ಪರಿಣಾಮವಾಗಿ, ಕರ್ನೆಲ್ ಅನ್ನು ಅಭಿವೃದ್ಧಿಪಡಿಸುವಾಗ, ಯಾವ ಗಣನೆಯನ್ನು ಮಾಡಬೇಕು ಎಂಬುದನ್ನು ವ್ಯಕ್ತಪಡಿಸುವುದರಲ್ಲಿ ಸಂಕೀರ್ಣತೆಯು ಸಾಮಾನ್ಯವಾಗಿ ಇರುವುದಿಲ್ಲ, ಆದರೆ ಹಾರ್ಡ್‌ವೇರ್ ಅದನ್ನು ನಿಖರವಾಗಿ ಹೇಗೆ ಮಾಡಬೇಕು ಎಂಬುದರಲ್ಲಿ ಇರುತ್ತದೆ. ಅತ್ಯುತ್ತಮ ಕರ್ನೆಲ್‌ಗಳನ್ನು ನಿರ್ಮಿಸಲು, ಆಧಾರವಾಗಿರುವ ಹಾರ್ಡ್‌ವೇರ್ ಮೇಲೆ ಸರಿಯಾದ ಅಮಸ್ಟ್ರ್ಯಾಕ್ಷನ್‌ಗಳನ್ನು ನೀಡುವ ಭಾಷೆಯ ನಮಗೆ ಅಗತ್ಯವಿದೆ. ಸರಿಯಾದ ಭಾಷೆಯು ಗಣನೆಯನ್ನು ಮಾಡಲು ಉನ್ನತ ಮಟ್ಟದ ಕಾರ್ಯಾಚರಣೆಗಳನ್ನು ನಿರ್ದಿಷ್ಟಪಡಿಸಲು ನಮಗೆ ಅನುವು ಮಾಡಿಕೊಡುತ್ತದೆ, ವಿಶಾಲವಾದ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್‌ನಲ್ಲಿರುವ ಎಲ್ಲಾ ಹಂತಗಳಲ್ಲಿ ಅವುಗಳನ್ನು ಆಕ್ರಮಣಕಾರಿಯಾಗಿ ವಿಶೇಷಗೊಳಿಸುತ್ತದೆ ಮತ್ತು ಆಪ್ಟಿಮೈಜ್ ಮಾಡುತ್ತದೆ.

CuTeDSL ಎನ್ನುವುದು CuTe ಲೇಔಟ್ ಆಲ್ಜೀಬ್ರಾ ಮತ್ತು MLIR ಮೇಲೆ ನಿರ್ಮಿಸಲಾದ NVIDIA GPU ಗಳಿಗಾಗಿ ಪೈಥಾನ್-ಆಧಾರಿತ DSL ಆಗಿದೆ. ಇದು ಕಡಿಮೆ-ಮಟ್ಟದ ಹಾರ್ಡ್‌ವೇರ್ ಪ್ರಿಮಿಟಿವ್‌ಗಳ ಮೇಲೆ ನಿಯಂತ್ರಣವನ್ನು ಉಳಿಸಿಕೊಂಡುಕೊಂಡು, ಹೆಚ್ಚು ಆಪ್ಟಿಮೈಸ್ ಮಾಡಿದ PTX ಗೆ ಜಸ್ಟ್-ಇನ್-ಟೈಮ್ ಕಂಪೈಲ್ ಮಾಡುತ್ತದೆ. ಇದು ಅಭಿವೃದ್ಧಿ ವೇಗಕ್ಕೆ ಅಡ್ಡಿಪಡಿಸುವ ದುಬಾರಿ ಕಂಪೈಲ್ ಸಮಯವಿಲ್ಲದೆ, CUTLASS C++ ನಂತಹದೇ ವೈಶಿಷ್ಟ್ಯಗಳು ಮತ್ತು ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ನೀಡುವ ಗುರಿಯನ್ನು ಹೊಂದಿದೆ. ನಾವು ಸಮೀಕ್ಷೆ ಮಾಡಿದ ಅಥವಾ ಕೆಲಸ ಮಾಡಿದ ಭಾಷೆಗಳಲ್ಲಿ (CUDA, cuTile, TVM IR ಮತ್ತು Triton ಸೇರಿದಂತೆ), ಇದು ನಮ್ಮ ಕರ್ನೆಲ್‌ಗಳ ಅಭಿವೃದ್ಧಿಗೆ ಅತ್ಯಂತ ಸೂಕ್ತವಾಗಿದೆ ಎಂದು ನಾವು ಪರಿಗಣಿಸುತ್ತೇವೆ.

ಒಟ್ಟಾರೆಯಾಗಿ, Triton ಅಥವಾ cuTile ನಷ್ಟು ಉನ್ನತ ಮಟ್ಟದ ಮತ್ತು ಸಂಕ್ಷಿಪ್ತವಾಗಿಲ್ಲದಿದ್ದರೂ, CuTe ಲೇಔಟ್ ಆಲ್ಜೀಬ್ರಾ ಮತ್ತು ಅದರ ಸಂಯೋಜನೆಗಳು (cute.Tensor ಮೂಲಕ ಬಹಿರಂಗಪಡಿಸಲಾಗಿದೆ) ಸಮಂಜಸವಾದ ಬದಲಾವಣೆಗಳಾಗಿವೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ. ಹೆಚ್ಚುವರಿಯಾಗಿ, CuTeDSL ಹಾರ್ಡ್‌ವೇರ್ ಪ್ರಿಮಿಟಿವ್‌ಗಳಿಗೆ ನಿರ್ಬಂಧವಿಲ್ಲದ ಪ್ರವೇಶವನ್ನು ನೀಡುತ್ತದೆ, ಎಲ್ಲಾ ಎಡ್ಜ್ ಕೇಸ್‌ಗಳಲ್ಲಿ ಗರಿಷ್ಠ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಸಾಧಿಸಲು ನಮಗೆ ಅಗತ್ಯವಿರುವ ನಿಯಂತ್ರಣ ಮಟ್ಟವನ್ನು ಒದಗಿಸುತ್ತದೆ. CUDA ನೊಂದಿಗೆ, ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಪೆಷಲೈಸೇಶನ್‌ಗೆ ಅವಕಾಶಗಳು ಮಿತವಾಗಿದ್ದವು, ಏಕೆಂದರೆ ಪ್ರತಿ ಹೆಚ್ಚುವರಿ ನಿಯರಾಚನೆಯು ಕಂಪೈಲ್ ಸಮಯಕ್ಕೆ ಅಸಮಂಜಸವಾದ ವೆಚ್ಚವನ್ನು ಅಥವಾ ಬಿಲ್ಡ್ ಸಿಸ್ಟಮ್‌ಗೆ ಸಂಕೀರ್ಣತೆಯನ್ನು ಸೇರಿಸಿತು. ಇದಕ್ಕೆ ವ್ಯತಿರಿಕ್ತವಾಗಿ, ಜಸ್ಟ್-ಇನ್-ಟೈಮ್ ಕಂಪೈಲ್ಡ್ DSL ನೊಂದಿಗೆ, ಪುನರಾವರ್ತನೆಯ ವೇಗವನ್ನು ತ್ಯಾಗ ಮಾಡದೆ ನಾವು ಹೆಚ್ಚು ಆಕ್ರಮಣಕಾರಿಯಾಗಿ ಸ್ಪೆಷಲೈಸ್ ಮಾಡಬಹುದು. CUDA ಅನ್ನು ಮಧ್ಯಂತರ ಹಂತವಾಗಿ ಬಳಸಿದ TVM IR ಗಿಂತ ಭಿನ್ನವಾಗಿ, PTX ಗೆ ಲೋವರ್ ಆದ MLIR-ಆಧಾರಿತ ಪೈಪ್‌ಲೈನ್ CUDA ಅನ್ನು ಮಧ್ಯಂತರ ಹಂತವಾಗಿ ಹೊರಸೂಸುವ ಪೈಪ್‌ಲೈನ್‌ಗಿಂತ ಹೆಚ್ಚು ದೃಢವಾದ ಮತ್ತು ಡೀಬಗ್ ಮಾಡಲು ಸುಲಭವಾಗಿದೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ. ಹೆಚ್ಚುವರಿಯಾಗಿ, PTX ಸೂಚನೆಗಳನ್ನು ಸುತ್ತುವರೆದಿರಲು ಕಸ್ಟಮ್ ಆಪ್‌ಗಳನ್ನು ಬಳಸುವುದರಿಂದ, ಆಧಾರವಾಗಿರುವ ಹಾರ್ಡ್‌ವೇರ್‌ನ ಎಲ್ಲಾ ವೈಶಿಷ್ಟ್ಯಗಳನ್ನು ನಾವು ಸಂಪೂರ್ಣವಾಗಿ ಬಳಸಿಕೊಳ್ಳಬಹುದು ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ.

CUDA ಕರ್ನೆಲ್‌ಗಳ ಕಂಪೈಲ್ ವೆಚ್ಚದ ಬಹುಪಾಲು ಭಾಗವು ಸಾಮಾನ್ಯವಾಗಿ ಕರ್ನೆಲ್‌ಗಳ ಟೆಂಪ್ಲೇಟ್ ಎಕ್ಸ್‌ಪ್ಯಾನ್ಷನ್ ಹಂತದಲ್ಲಿ ಕಳೆಯಲಾಗುತ್ತದೆ, ಏಕೆಂದರೆ PTX ಹಂತದಿಂದ ಮುಂದಕ್ಕೆ ಆಪ್ಟಿಮೈಸೇಶನ್ ಮತ್ತು ಕಂಪೈಲ್ ಅನ್ನು ಕೆಲವೇ ಸೆಕೆಂಡುಗಳಲ್ಲಿ ಪೂರ್ಣಗೊಳಿಸಬಹುದು. ಅನೇಕ ಕಂಪೈಲ್-ಸಮಯದ ಪ್ಯಾರಾಮೀಟರ್‌ಗಳೊಂದಿಗೆ, ನಾನ್‌-ಟ್ರಿವಿಯಲ್ ಕರ್ನೆಲ್‌ಗಾಗಿ ಕಂಪೈಲ್‌ನ ಈ ಹಂತವು ಬಹು ನಿಮಿಷಗಳವರೆಗೆ ಸ್ಫೋಟಗೊಳ್ಳಬಹುದು. ಸಮಾಂತರ ಕಂಪೈಲ್‌ಗೆ ಅವಕಾಶಗಳು ನಂತರದಲ್ಲಿ, ಅಂದರೆ IR ಮತ್ತು PTX ಮಟ್ಟಗಳಲ್ಲಿ ಮಾತ್ರ ಉದ್ಭವಿಸುತ್ತವೆ ಎಂಬ ಅಂಶದಿಂದ ಈ ಬಾಟಲ್‌ನೆಕ್ ಒತ್ತಿಹೇಳಲ್ಪಡುತ್ತದೆ, ಜೊತೆಗೆ C++ ಟೆಂಪ್ಲೇಟ್ ನಿರ್ವಹಣೆಯು ಅನುಕ್ರಮವಾಗಿ ಸಂಭವಿಸುತ್ತದೆ. ಇದನ್ನು ಪರಿಹರಿಸಲು ತಂತ್ರಗಳು ಅಸ್ತಿತ್ವದಲ್ಲಿವೆಯಾದರೂ, ಟೆಂಪ್ಲೇಟ್ ಇನ್‌ಸ್ಟಾನ್ಸಿಯೇಷನ್‌ಗಳನ್ನು ಪ್ರತ್ಯೇಕ ಕಂಪೈಲ್ ಯುನಿಟ್‌ಗಳಿಗೆ ಸರಿಸುವ ಮೂಲಕ, ಅವುಗಳಿಗೆ ಬಿಲ್ಡ್ ಸಿಸ್ಟಮ್‌ಗಳಿಗೆ ಗಮನಾರ್ಹವಾದ ಟ್ವೀಕ್‌ಗಳು ಮತ್ತು ಬಳಕೆಗೆ ಕಷ್ಟಕರವಾದ ಮ್ಯಾಕ್ರೋಗಳ ಅಗತ್ಯವಿರುತ್ತದೆ. ಇದಕ್ಕೆ ವ್ಯತಿರಿಕ್ತವಾಗಿ, ಜಸ್ಟ್-ಇನ್-ಟೈಮ್ (JIT) ಕಂಪೈಲ್ ಮಾಡಲಾದ CuTeDSL ಕರ್ನೆಲ್‌ಗಳಿಗೆ ಸಂಕೀರ್ಣ ಬಿಲ್ಡ್ ನಿಯಮಗಳ ಅಗತ್ಯವಿರುವುದಿಲ್ಲ ಮತ್ತು ಕಂಪೈಲ್ ಸಮಯಕ್ಕೆ ಅಡ್ಡಿಪಡಿಸುವುದಿಲ್ಲ, ಏಕೆಂದರೆ ಅಭಿವೃದ್ಧಿ ಚಕ್ರವು ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್‌ನಲ್ಲಿ ಒಂದೇ ಬಿಂದುವಿನ ಮೇಲೆ ಕೇಂದ್ರೀಕರಿಸಬಹುದು, ಆದರೆ ನಂತರ ಎಚ್ಚರಿಕೆಯ ಯುನಿಟ್ ಟೆಸ್ಟ್‌ಗಳು ಕರ್ನೆಲ್‌ನ ಸಂಪೂರ್ಣ ವಿನ್ಯಾಸ ಸ್ಪೇಸ್ ಅನ್ನು ಒಳಗೊಳ್ಳಬಹುದು.

JIT ಕಂಪೈಲ್ ವೆಚ್ಚವನ್ನು ಅಭಿವೃದ್ಧಿ-ಸಮಯದಿಂದ ರನ್‌ಟೈಮ್‌ಗೆ ಸರಿಸಿದರೂ, ಓವರ್‌ಹೆಡ್ ಸ್ವೀಕಾರಾರ್ಹಕ್ಕಿಂತ ಹೆಚ್ಚು. ಮಾದರಿಯನ್ನು ನಿಯೋಜಿಸಿದಾಗ, ಇದು ಸಂಪೂರ್ಣ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್‌ನಿಂದ ಕೆಲವೇ ಪಾಯಿಂಟ್‌ಗಳನ್ನು ಸಕ್ರಿಯಗೊಳಿಸುತ್ತದೆ, ಇನ್‌ಫೆರೆನ್ಸ್ ಎಂಜಿನ್‌ನ ಸ್ಟಾರ್ಟ್‌ಅಪ್/ವಾರ್ಮ್ಅಪ್ ಸಮಯದಲ್ಲಿ ಸೆಕೆಂಡ್-ಪ್ರಮಾಣದ ಓವರ್‌ಹೆಡ್‌ಗಳನ್ನು ಉಂಟುಮಾಡುತ್ತದೆ. ಇದಕ್ಕೆ ವ್ಯತಿರಿಕ್ತವಾಗಿ, ಕರ್ನೆಲ್‌ನ ಎಲ್ಲಾ ಕಾನ್ಫಿಗರೇಶನ್‌ಗಳನ್ನು ಪ್ರೀ-ಕಂಪೈಲ್ ಮಾಡುವುದು ಬಹು ನಿಮಿಷಗಳವರೆಗೆ ವಿಸ್ತರಿಸಬಹುದು. ಹೆಚ್ಚುವರಿಯಾಗಿ, JIT-ಕಂಪೈಲ್ ಮಾಡಿದ ಆರ್ಟಿಫ್ಯಾಕ್ಟ್ಸ್‌ಗಳನ್ನು ಸಂಗ್ರಹಿಸುವ ಮೂಲಕ ಓವರ್‌ಹೆಡ್‌ಗಳನ್ನು ಮತ್ತಷ್ಟು ತಗ್ಗಿಸಬಹುದು.

CuTe ಮತ್ತು CUTLASS ನಂತೆ ಟೆಂಪ್ಲೇಟ್-ಆಧಾರಿತ ಡಿಸ್ಪ್ಯಾಚ್ ಮತ್ತು ಸ್ಪೆಷಲೈಸೇಶನ್ ಅನ್ನು ಬಳಸಿದಾಗ ದೋಷಗಳ ಮೂಲವನ್ನು ಟ್ರೇಸ್ ಮಾಡುವ ಕಷ್ಟದಿಂದ CUDA ಯ ಭಾರವಾದ ಕಂಪೈಲ್ ವೆಚ್ಚವು ಒತ್ತಿಹೇಳಲ್ಪಡುತ್ತದೆ. DSL ನೊಂದಿಗೆ, ಅದನ್ನು ಉತ್ಪಾದಿಸಿದ ಕೋಡ್ ಸಾಲಿಗೆ ದೋಷಗಳನ್ನು ಸೂಚಿಸಬಹುದು, ಇದು ಡೀಬಗ್ ಮಾಡುವುದನ್ನು ಗಮನಾರ್ಹವಾಗಿ ಸುಲಭಗೊಳಿಸುತ್ತದೆ. ಹೆಚ್ಚುವರಿಯಾಗಿ, ಮಧ್ಯಂತರ MLIR ಹಂತಗಳ ಮೂಲಕ ಅಥವಾ ಉತ್ಪತ್ತಿಯಾದ PTX ಗೆ ಲಗತ್ತಿಸಲಾದ ಡೀಬಗ್ ಮಾಡುವ ಮಾಹಿತಿಯನ್ನು ಅವಲಂಬಿಸಿ ಉತ್ಪತ್ತಿಯಾದ ಕೋಡ್ ಅನ್ನು ಅದನ್ನು ಹುಟ್ಟುಹಾಕಿದ DSL ಸ್ಟೇಟ್‌ಮೆಂಟ್‌ಗೆ ಮ್ಯಾಪ್ ಮಾಡುವುದು ತುಂಬಾ ಸರಳವಾಗಿದೆ. DSL ನ ಈ ಅಂಶಗಳಿಗೆ ಧನ್ಯವಾದಗಳು, ಕರ್ನೆಲ್‌ಗಳನ್ನು ಅಭಿವೃದ್ಧಿಪಡಿಸುವಾಗ ನಾವು ಹೆಚ್ಚು ಉತ್ಪಾದಕರಾಗಿದ್ದೇವೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ.

ಕರ್ನೆಲ್‌ಗಳು ಮಾದರಿಗಳಾದ್ಯಂತ ಅನೇಕ ಸಾಮ್ಯತೆಗಳನ್ನು ಹಂಚಿಕೊಂಡರೂ, ಅವು ಇನ್ನೂ ಕಾರ್ಯಕ್ಷಮತೆಯ ಮೇ I ಬೀರುವ ಶ್ರೇಣಿಯ ಕಾನ್ಫಿಗರೇಶನ್ ನಾಬ್‌ಗಳೊಂದಿಗೆ ಸಜ್ಜುಗೊಂಡಿರಬೇಕು. ಸಂಬಂಧಿತ ಪ್ಯಾರಾಮೀಟರ್‌ಗಳು ಇವುಗಳನ್ನು ಒಳಗೊಂಡಿವೆ:

ಹಿಡನ್ ಡೈಮೆನ್ಷನ್/ಹೆಡ್ ಡೈಮೆನ್ಷನ್/ಅಟೆನ್ಷನ್ ಹೆಡ್‌ಗಳ ಸಂಖ್ಯೆ

ಇನ್‌ಪುಟ್/ಔಟ್‌ಪುಟ್ ಡೇಟಾ ಪ್ರಕಾರಗಳು

ಕ್ವಾಂಟೈಸೇಶನ್ ಯೋಜನೆ

ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಥಿರಾಂಕಗಳು (ಬಯಾಸ್‌ಗಳು, ಆಕ್ಟಿವೇಷನ್ ಕಾರ್ಯಗಳು ಇತ್ಯಾದಿ)

ROSE ಆರ್ಕಿಟೆಕ್ಚರ್, ಅದರ ಎಂಜಿನ್, ಮಾದರಿ, ಕೆರ್ನಲ್ ಲೈಬ್ರರಿ ಮತ್ತು ಸೇವೆಗಳನ್ನು ಒಳಗೊಂಡಂತೆ.

ಕಾರ್ಯಕ್ಷಮತೆ ಮತ್ತು ಡೀಬಗ್ ಮಾಡುವಿಕೆ ಎರಡಕ್ಕೂ ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಪೆಷಲೈಸೇಶನ್ ಉಪಯುಕ್ತವಾಗಿದೆ. ವಿಶೇಷವಾಗಿ ಪೈಪ್‌ಲೈನ್ ಮಾಡಿದ ಪ್ರೊಸೆಸರ್‌ನಲ್ಲಿ, GPU ಗಳ ವಿಷಯದಲ್ಲಿರುವಂತೆ, ಯಾವಾಗಲೂ ತೆಗೆದುಕೊಳ್ಳುವ ಅಥವಾ ಎಂದಿಗೂ ತೆಗೆದುಕೊಳ್ಳದ ಬ್ರಾಂಚ್‌ಗಿಂತ ಅಸ್ತಿತ್ವದಲ್ಲಿಲ್ಲದ ಬ್ರಾಂಚ್ ಉತ್ತಮವಾಗಿರುತ್ತದೆ. ಹೆಚ್ಚುವರಿಯಾಗಿ, ನಿರ್ದಿಷ್ಟ ಕಾನ್ಫಿಗರೇಶನ್ ನಾಬ್‌ಗಳ ಸೆಟ್‌ಗಾಗಿ ಕರ್ನೆಲ್ ಅನ್ನು ಕಂಪೈಲ್ ಮಾಡುವುದರಿಂದ ಕಡಿಮೆ PTX/SASS ಸೂಚನೆಗಳು ದೊರೆಯುತ್ತವೆ, ತನಿಖೆಯ ಸಮಯದಲ್ಲಿ cuda-gdb ಅನ್ನು ಬಳಸುವಾಗ ಇದನ್ನು ವಿಶ್ಲೇಷಿಸಲು ಮತ್ತು ಡೀಬಗ್ ಮಾಡಲು ಸುಲಭವಾಗುತ್ತದೆ.

ನಿರ್ದಿಷ್ಟ ಬ್ಯಾಚ್ ಗಾತ್ರ ಮತ್ತು ಕಾನ್ಫಿಗರೇಶನ್‌ಗಾಗಿ, DRAM ಗೆ ನಿರೀಕ್ಷಿತ ಟ್ರಾಫಿಕ್ ಮತ್ತು ಅಂಕಗಣಿತದ ಕಾರ್ಯಾಚರಣೆಗಳ ಸಂಖ್ಯೆಯನ್ನು ಅಂದಾಜು ಮಾಡಬಹುದು. ತರುವಾಯ, ಅತ್ಯುತ್ತಮ ಗ್ರಿಡ್ ಹಂಚಿಕೆ, ಸಿಂಕ್ರೊನೈಸೇಶನ್ ಮತ್ತು ಕೆಲವು ಹಾರ್ಡ್‌ವೇರ್ ವೈಶಿಷ್ಟ್ಯಗಳ ಬಳಕೆಗಾಗಿ ಪ್ರತಿ ಕೇಸ್-ಬೈ-ಕೇಸ್ ಆಧಾರದ ಮೇಲೆ ಅತ್ಯುತ್ತಮ ನಿರ್ಧಾರಗಳನ್ನು ತೆಗೆದುಕೊಳ್ಳಬಹುದು. ಸ್ವೀಕಾರಾರ್ಹವಲ್ಲದ ಕಂಪೈಲ್-ಸಮಯದ ಓವರ್‌ಹೆಡ್ ಅನ್ನು ಉಂಟುಮಾಡದೆ ಆಳವಾದ ಸ್ಪೆಷಲೈಸೇಶನ್‌ಗಾಗಿ ಘಟಕಗಳನ್ನು ವ್ಯಾಖ್ಯಾನಿಸಲು ಮತ್ತು ಮರುಬಳಕೆ ಮಾಡಲು CuTeDSL ನಮಗೆ ಅನುಮತಿಸುತ್ತದೆ.

ಸಾಮಾನ್ಯವಾಗಿ, ಇನ್‌ಫೆರೆನ್ಸ್ ಕರ್ನೆಲ್‌ಗಳ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಎರಡು ಪ್ರತ್ಯೇಕ ಬಳಕೆ ಪ್ರಕರಣಗಳಿಗೆ ಆಪ್ಟಿಮೈಜ್ ಮಾಡಬೇಕು:

ಪ್ರಿಫಿಲ್: ಸಣ್ಣ ಬ್ಯಾಚ್ ಗಾತ್ರ (ಕೆಲವು ಚಂಕ್‌ಗಳು, ಅನುಕ್ರಮಗಳು), ಪ್ರತಿ ಅನುಕ್ರಮಕ್ಕೆ ಅನೇಕ ಟೋಕನ್‌ಗಳು. ಎಂಭೆಡ್ಡಿಂಗ್‌ಗಳು ಮತ್ತು LLM ಗಳ ಇನ್‌ಪುಟ್ ಟೋಕನ್‌ಗಳೆರಡಕ್ಕೂ ಥ್ರುಪುಟ್ ಅನ್ನು ಗರಿಷ್ಠಗೊಳಿಸಬೇಕು. ಹತ್ತರಿಂದ ನೂರಾರು ಮಿಲಿಸೆಕೆಂಡ್‌ಗಳಲ್ಲಿ ಲೇಟೆನ್ಸಿ ಇರುತ್ತದೆ.

ಡಿಕೋಡ್: ಸಣ್ಣದರಿಂದ ದೊಡ್ಡ ಬ್ಯಾಚ್ ಗಾತ್ರಗಳು, ಪ್ರತಿ ಅನುಕ್ರಮಕ್ಕೆ ಒಂದು (ಊಹಾತ್ಮಕ ಡಿಕೋಡಿಂಗ್ ಇಲ್ಲ) ರಿಂದ ಕೆಲವು ಟೋಕನ್‌ಗಳು (ಡ್ರಾಫ್ಟ್-ಟಾರ್ಗೆಟ್ ಸ್ಪೆಕ್ಯುಲೇಷನ್, ಮಲ್ಟಿ-ಟೋಕನ್ ಪ್ರಿಡಿಕ್ಷನ್). ಕರ್ನೆಲ್ ಲಾಂಚ್ ಲೇಟೆನ್ಸಿಗೆ ಹತ್ತಿರವಿರುವ ಲೇಟೆನ್ಸಿ.

GEMM ಮತ್ತು attentionಗಾಗಿ, prefill ಹಂತವು ಹೆಚ್ಚಿನ ಸಂಖ್ಯೆಯ ಟೋಕನ್‌ಗಳ ಮೇಲೆ ನಡೆಯುತ್ತದೆ ಮತ್ತು ಕಾರ್ಯಕ್ಷಮತೆ ಸಾಧನದ compute throughput‌ನಿಂದ ನಿರ್ಬಂಧಿತವಾಗಿರುತ್ತದೆ. decoding ನಲ್ಲಿ ಸಣ್ಣ ಬ್ಯಾಚ್ ಗಾತ್ರಗಳ ಮೇಲೆ ಅದೇ ಕಾರ್ಯಾಚರಣೆ ನಡೆಯುತ್ತದೆ; ಇದರ ನಿರ್ವಹಣಾ ಸಮಯಗಳು DRAM‌ನಿಂದ streaming multiprocessors ಮತ್ತು tensor cores‌ಗೆ activations ಮತ್ತು weights ಅನ್ನು ಲೋಡ್ ಮಾಡುವ ವೆಚ್ಚದಿಂದ ಆವರಿತವಾಗುತ್ತವೆ, ಮತ್ತು ಬದಲು ಸಾಧನದ memory bandwidth‌ನಿಂದ ನಿರ್ಬಂಧಿತವಾಗಿರುತ್ತವೆ. ಬ್ಯಾಚ್ ಗಾತ್ರ ಅಥವಾ ಬ್ಯಾಚ್‌ನಲ್ಲಿನ ಟೋಕನ್‌ಗಳ ಒಟ್ಟು ಸಂಖ್ಯೆ ತುಲನಾತ್ಮಕವಾಗಿ ಸ್ಪಷ್ಟ ಸೂಚಕವಾಗಿದ್ದರೂ, prefill ಮತ್ತು decode ನಡುವಿನ switchover threshold ಹೆಚ್ಚು ಸೂಕ್ಷ್ಮವಾಗಿದೆ, ಏಕೆಂದರೆ memory bandwidth ಕೂಡ ಮಾದರಿಯ hidden dimension ಹಾಗೂ activation‌ನ data type‌ನಿಂದ ನಿರ್ಧಾರಗೊಳ್ಳುತ್ತದೆ. ಆದ್ದರಿಂದ, algorithmic ವ್ಯತ್ಯಾಸಗಳು ಅಥವಾ grid/block ನಿಯೋಜನೆಗಳು ಯಾವುದಾದರೂ ಒಂದಕ್ಕೆ ಲಾಭಕರವಾಗಿದ್ದರೆ kernelsಗಳನ್ನು prefill ಮತ್ತು decode implementationsಗಳಾಗಿ ವಿಭಜಿಸಲಾಗುತ್ತದೆ.

ಪ್ರಿಫಿಲ್ ಮತ್ತು ಡಿಕೋಡ್ ಕರ್ನೆಲ್‌ಗಳ ನಡುವೆ ಸಾರ್ವತ್ರಿಕವಾಗಿ ಭಿನ್ನವಾಗಿರುವ ಪ್ಯಾರಾಮೀಟರ್ ಎಂದರೆ ಬ್ಲಾಕ್‌ನಲ್ಲಿರುವ ವಾರ್ಪ್‌ಗಳ ಸಂಖ್ಯೆ (ಥ್ರೆಡ್‌ಗಳ ಸಂಖ್ಯೆ, ಪ್ರತಿ ವಾರ್ಪ್‌ಗೆ 32 ಥ್ರೆಡ್‌ಗಳ ವಾರ್ಪ್ ಗಾತ್ರದಿಂದ ಭಾಗಿಸಲಾಗಿದೆ). ರಿಜಿಸ್ಟರ್ ಹಂಚಿಕೆಯಲ್ಲಿ ಇದು ಸಹಾಯ ಮಾಡುವುದರಿಂದ, ಈ ಪ್ಯಾರಾಮೀಟರ್ ಅನ್ನು ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಥಿರಾಂಕವನ್ನಾಗಿ ಮಾಡುವುದು ಅಪೇಕ್ಷಣೀಯವಾಗಿದೆ. CUDA ನಲ್ಲಿ, __launch_bounds__ ಮೂಲಕ ಥ್ರೆಡ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ನಿರ್ದಿಷ್ಟಪಡಿಸುವ ಮೂಲಕ ಮತ್ತು ವಾರ್ಪ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ಟೆಂಪ್ಲೇಟ್ ಪ್ಯಾರಾಮೀಟರ್‌ನಂತೆ ಒದಗಿಸುವ ಮೂಲಕ ಇದನ್ನು ಸಾಮಾನ್ಯವಾಗಿ ಮಾಡಲಾಗುತ್ತದೆ:

cpp

template <size_t NUM_WARPS>
__global__ void 
__launch_bounds__(NUM_WARPS * WARP_SIZE, 1)
some_kernel(...) {

}

ಬ್ಲಾಕ್‌ನಲ್ಲಿರುವ ಎಲ್ಲಾ ಥ್ರೆಡ್‌ಗಳಿಗೆ ಲಭ್ಯವಿರುವ ರಿಜಿಸ್ಟರ್‌ಗಳ ಸಂಖ್ಯೆಯು ಹಾರ್ಡ್‌ವೇರ್ ನಿರ್ಬಂಧವಾಗಿದೆ, ಇದು Hopper ಮತ್ತು Blackwell ನಲ್ಲಿ 64K ಆಗಿದೆ. ಥ್ರೆಡ್‌ಗಳ ಸಂಖ್ಯೆಯು ಮುಂಚಿತವಾಗಿ ತಿಳಿದಿದ್ದರೆ, ಪ್ರತಿ ಥ್ರೆಡ್‌ಗೆ ಎಷ್ಟು ರಿಜಿಸ್ಟರ್‌ಗಳನ್ನು ನಿಯೋಜಿಸಬಹುದು ಎಂಬುದನ್ನು ಕಂಪೈಲರ್ ನಿಖರವಾಗಿ ತಿಳಿದಿರುತ್ತದೆ ಮತ್ತು ಹೆಚ್ಚು ಆಕ್ರಮಣಕಾರಿಯಾಗಿ ಆಪ್ಟಿಮೈಸ್ ಮಾಡಿದ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಕಂಪೈಲ್ ಮಾಡುವಲ್ಲಿ ಯಶಸ್ವಿಯಾಗಬಹುದು. ಆದಾಗ್ಯೂ, ಇದು ಕಂಪೈಲ್-ಸಮಯದ ವೆಚ್ಚದಲ್ಲಿ ಬರುತ್ತದೆ, ಏಕೆಂದರೆ ಈಗ ಕರ್ನೆಲ್ ಅನ್ನು ವಿವಿಧ ವಾರ್ಪ್ ಗಾತ್ರಗಳಿಗಾಗಿ ಮರುಕಂಪೈಲ್ ಮಾಡಬೇಕಾಗುತ್ತದೆ. CuTeDSL ನೊಂದಿಗೆ, ಇದು ಇನ್ನು ಮುಂದೆ ಸಮಸ್ಯೆಯಲ್ಲ, ಏಕೆಂದರೆ ಸಂಪೂರ್ಣ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್‌ಗೆ ನಾವು ಮುಂಚಿತವಾಗಿ ಕಂಪೈಲ್ ವೆಚ್ಚವನ್ನು ಪಾವತಿಸಬೇಕಾಗಿಲ್ಲ. CUDA ನಲ್ಲಿ __latch_bounds__ ನಿಂದ ಸೂಚಿಸಲಾದ ಆಪ್ಟಿಮೈಸೇಶನ್ ಅನ್ನು ಸೂಚ್ಯವಾಗಿ ನಿರ್ವಹಿಸುವ ಮೂಲಕ, ಬ್ಲಾಕ್ ಗಾತ್ರವು ಸ್ಥಿರ ಅಭಿವ್ಯಕ್ತಿಯಾಗಿರಬೇಕು ಎಂದು DSL ನಿರೀಕ್ಷಿಸುತ್ತದೆ.

ವಾರ್ಪ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ಸ್ಥಿರಾಂಕವನ್ನಾಗಿ ಮಾಡುವುದರ ಹೊರತಾಗಿ, ಪ್ರಿಫಿಲ್ ಮತ್ತು ಡಿಕೋಡ್ ನಡುವೆ ಅದನ್ನು ಬದಲಾಯಿಸುವುದು ಅಪೇಕ್ಷಣೀಯವಾಗಿದೆ. GPU ನ ಮೆಮೊರಿ ಬ್ಯಾಂಡ್‌ವಿಡ್ತ್ ಅನ್ನು ಸಂಪೂರ್ಣವಾಗಿ ಬಳಸಿಕೊಳ್ಳಲು, ಪ್ರತಿ ಬ್ಲಾಕ್‌ಗೆ ಹೆಚ್ಚಿನ ಸಂಖ್ಯೆಯ ವಾರ್ಪ್‌ಗಳನ್ನು ಯಾವಾಗಲೂ ನಿಯೋಜಿಸುವುದು ಸಾಕಾಗುವುದಿಲ್ಲ, ಏಕೆಂದರೆ ಲಭ್ಯವಿರುವ ಸ್ಟ್ರೀಮಿಂಗ್ ಮಲ್ಟಿಪ್ರೊಸೆಸರ್‌ಗಳಲ್ಲಿ (SM ಗಳು) ಹೆಚ್ಚಿನವುಗಳನ್ನು ಬಳಸುವ ಮೂಲಕ ಮಾತ್ರ ಸಂಪೂರ್ಣ ಬ್ಯಾಂಡ್‌ವಿಡ್ತ್ ಅನ್ನು ಸಾಧಿಸಬಹುದು. ಪರಿಣಾಮವಾಗಿ, ನಿರ್ದಿಷ್ಟ ಸಂಖ್ಯೆಯ ಟೋಕನ್‌ಗಳು ಮತ್ತು ತಿಳಿದಿರುವ ಡೇಟಾ ಪ್ರಕಾರಗಳು ಮತ್ತು ಹಿಡನ್ ಡೈಮೆನ್ಶನ್‌ಗಳಿಗಾಗಿ, ಉತ್ಪತ್ತಿಯಾದ ಕೋಡ್ 128-ಬಿಟ್ ಲೋಡ್/ಸ್ಟೋರ್ ಸೂಚನೆಗಳನ್ನು (Blackwell ನಲ್ಲಿ 256-ಬಿಟ್) ಬಳಸುವುದನ್ನು ಖಚಿತಪಡಿಸಿಕೊಳ್ಳಲು ನಾವು ಮೊದಲು ಆದರ್ಶವಾದ ಪ್ರತಿ-ಥ್ರೆಡ್ ವೆಕ್ಟರ್ ಡೈಮೆನ್ಶನ್ ಅನ್ನು ಆರಿಸಿಕೊಳ್ಳುತ್ತೇವೆ (bfloat16 ಗೆ 8, float32 ಗೆ 4). ವೆಕ್ಟರ್ ಗಾತ್ರವನ್ನು ತಿಳಿದುಕೊಂಡು, ಕರ್ನೆಲ್‌ಗೆ ಅಗತ್ಯವಿರುವ ವಾರ್ಪ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ನಾವು ಲೆಕ್ಕ ಹಾಕಬಹುದು. ನಾವು ಮೊದಲು SM ಗಳ ಸಂಖ್ಯೆಯನ್ನು ಗರಿಷ್ಠಗೊಳಿಸಲು ಪ್ರಯತ್ನಿಸುತ್ತೇವೆ: SM ಗಳಿಗಿಂತ ಕಡಿಮೆ ವಾರ್ಪ್‌ಗಳಿದ್ದರೆ, ಪ್ರತಿ ಬ್ಲಾಕ್‌ಗೆ ಒಂದೇ ವಾರ್ಪ್ ಅನ್ನು ಆಯ್ಕೆ ಮಾಡಲಾಗುತ್ತದೆ. ಇದು ಸಾಮಾನ್ಯವಾಗಿ ಕರ್ನೆಲ್‌ನ ಡಿಕೋಡ್ ರೂಪಾಂತರವಾಗಿದೆ. ಸಾಕಷ್ಟು ವಾರ್ಪ್‌ಗಳಿದ್ದರೆ, ನಾವು ಕ್ರಮೇಣ ವಾರ್ಪ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು 8, 16 ಅಥವಾ 32 ಕ್ಕೆ ಹೆಚ್ಚಿಸುತ್ತೇವೆ, ಮೇಲಿನ ಮಿತಿಯನ್ನು ಬೆಂಚ್‌ಮಾರ್ಕ್ ಮಾಡುತ್ತೇವೆ, ಅದರ ನಂತರ ನಾವು ಸಾಮಾನ್ಯವಾಗಿ ಕಾರ್ಯಕ್ಷಮತೆಯಲ್ಲಿ ಸುಧಾರಣೆಯನ್ನು ನೋಡುವುದಿಲ್ಲ.

Triton ಈ ಕೆಲವು ಆಪ್ಟಿಮೈಸೇಶನ್‌ಗಳನ್ನು ಹೂಡ್ ಅಡಿಯಲ್ಲಿ ನಿರ್ವಹಿಸಿದರೂ, ಉನ್ನತ ಮಟ್ಟದ ಕಂಪೈಲರ್‌ನ ನ್ಯೂನತೆಗಳಿಂದ ನಿರ್ಬಂಧಿಸಲ್ಪಡದೆ, CuTe ಲೇಔಟ್ ಆಲ್ಜೀಬ್ರಾವನ್ನು ಎಚ್ಚರಿಕೆಯಿಂದ ಬಳಸುವ ಮೂಲಕ ನಾವು ಒಂದೇ ರೀತಿಯ ಫಲಿತಾಂಶಗಳನ್ನು ಸಾಧಿಸಬಹುದು ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ. ವಿಶೇಷವಾಗಿ ಸಬ್-ವಾರ್ಪ್ ಸಮತಾದ ಕಡಿತಗಳು ಒಳಗೊಂಡಿರುವಾಗ, ಸೂಕ್ತವಾದ ಕೋಡ್ ಅನ್ನು ಉತ್ಪಾದಿಸಲು ಅಗತ್ಯವಿರುವ ಸೂಕ್ಷ್ಮ ನಿಯಂತ್ರಣವನ್ನು CuTeDSL ನಮಗೆ ನೀಡಿದೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ.

ಕೆಲವು ಸಂದರ್ಭಗಳಲ್ಲಿ, ಸಾಧ್ಯವಾದಷ್ಟು ಬ್ಲಾಕ್‌ಗಳಾದ್ಯಂತ ಕೆಲಸವನ್ನು ವಿತರಿಸಲು, ಎಲ್ಲಾ ಭಾಗವಹಿಸುವ ಬ್ಲಾಕ್‌ಗಳನ್ನು ಸಿಂಕ್ರೊನೈಸ್ ಮಾಡಲು cooperative_groups::this_grid().sync() ಅನ್ನು ಬಳಸಬಹುದು. ನಿರ್ಬಂಧಗಳು ಅನ್ವಯಿಸುತ್ತವೆ: ಬ್ಲಾಕ್ ಒಮ್ಮೆ ಕಾರ್ಯಗತಗೊಳ್ಳಲು ಪ್ರಾರಂಭಿಸಿದರೆ ಅದನ್ನು SM ನಿಂದ ಡಿ-ಶೆಡ್ಯೂಲ್ ಮಾಡಲು ಸಾಧ್ಯವಿಲ್ಲವಾದ್ದರಿಂದ, ಮುಂದಿನ ಪ್ರಗತಿಯನ್ನು ಖಾತರಿಪಡಿಸಲು, ಯಾವುದೇ ಹಂತದಲ್ಲಿ ಎಲ್ಲಾ ಬ್ಲಾಕ್‌ಗಳು ಸಿಂಕ್ರೊನೈಸೇಶನ್‌ನಲ್ಲಿ involucrated ಇರುವ ಗ್ಲೋಬಲ್ ಮೆಮೊರಿ ಸ್ಥಳಗಳನ್ನು ನವೀಕರಿಸಬಹುದು ಮತ್ತು ಪೋಲ್ ಮಾಡಬಹುದು ಎಂದು ಖಚಿತಪಡಿಸಿಕೊಳ್ಳಲು, ಬ್ಲಾಕ್‌ಗಳ ಸಂಖ್ಯೆಯು ಸಾಧನದಲ್ಲಿ ಲಭ್ಯವಿರುವ SM ಗಳ ಸಂಖ್ಯೆಯನ್ನು ಮೀರುವಂತಿಲ್ಲ. ಈ ನಿರ್ಬಂಧವನ್ನು cudaLaunchCooperativeKernel ನಿಂದ ಜಾರಿಗೆ ತರಲಾಗುತ್ತದೆ. Hopper ನಲ್ಲಿ ಗ್ರಿಡ್ ತಡೆಗೋಡೆಯ ವೆಚ್ಚವು ಸುಮಾರು 3us ಆಗಿದೆ, ಜೊತೆಗೆ ಕೆಲವು ಸಂದರ್ಭಗಳಲ್ಲಿ ಸೂಚ್ಯ ಮೆಮೊರಿ ತಡೆಗೋಡೆಯ ಮೂಲಕ ಎಲ್ಲಾ ಹಿಂದಿನ ಮೆಮೊರಿ ಕಾರ್ಯಾಚರಣೆಗಳನ್ನು ಫ್ಲಶ್ ಮಾಡುವ ವೆಚ್ಚವಿದೆ. ಒಟ್ಟಾರೆಯಾಗಿ, ಕರ್ನೆಲ್ ಅನ್ನು ಎರಡು ಲಾಂಚ್‌ಗಳಾಗಿ ವಿಭಜಿಸುವ ಓವರ್‌ಹೆಡ್‌ಗೆ ಇದು ಹೋಲಿಸಬಹುದಾಗಿದೆ, ಇದು ಕಾರ್ಯಕ್ಷಮತೆ-ಸಂವೇದನಾಶೀಲ ವಿನ್ಯಾಸ ನಿರ್ಧಾರಕ್ಕೆ ಕಾರಣವಾಗುತ್ತದೆ. ಪ್ರೋಗ್ರಾಮ್ಯಾಟಿಕ್ ಡಿಪೆಂಡೆಂಟ್ ಲಾಂಚ್ (PDL) Hopper ನಿಂದ ಮುಂದಕ್ಕೆ ಲಾಂಚ್ ಓವರ್‌ಹೆಡ್‌ಗಳನ್ನು ಕಡಿಮೆ ಮಾಡಬಹುದಾದರೂ, ಡಿಕೋಡ್ ಕರ್ನೆಲ್‌ಗಳಿಗೆ, ಓವರ್‌ಹೆಡ್ ಅನ್ನು ತಪ್ಪಿಸುವುದು ಯೋಗ್ಯವಾಗಿದೆ.

ಕೆಲಸದ ಹೊರೆಗಳಲ್ಲಿ ಒಂದಕ್ಕೆ SM ಗಳ ಸಂಖ್ಯೆಗಿಂತ ಹೆಚ್ಚು ಬ್ಲಾಕ್‌ಗಳ ಅಗತ್ಯವಿದ್ದರೆ, ಕರ್ನೆಲ್ ಲಾಂಚ್‌ಗಳ ಸೂಚ್ಯ ತಡೆಗೋಡೆಯನ್ನು ಅವಲಂಬಿಸಿ ಕರ್ನೆಲ್ ಅನ್ನು ವಿಂಗಡಿಸಬೇಕು. ಇದು ಪ್ರಾಥಮಿಕವಾಗಿ ಪ್ರಿಫಿಲ್‌ಗೆ ಅನ್ವಯಿಸುತ್ತದೆ. ಎಲ್ಲಾ ಕೆಲಸದ ಹೊರೆಗಳಿಗೆ ಆದರ್ಶವಾಗಿರುವ ಬ್ಲಾಕ್‌ಗಳ ಸಂಖ್ಯೆಯು ಗ್ರಿಡ್‌ಗೆ ಹೊಂದಿಕೊಳ್ಳುತ್ತಿದ್ದರೆ, ಗ್ರಿಡ್ ತಡೆಗೋಡೆಯನ್ನು ಆದ್ಯತೆ ನೀಡಲಾಗುತ್ತದೆ, ಕಡಿಮೆ ಸಂಖ್ಯೆಯ ಟೋಕನ್‌ಗಳ ಮೇಲೆ ಕಾರ್ಯನಿರ್ವಹಿಸುವ ಡಿಕೋಡ್ ಕರ್ನೆಲ್‌ಗಳ ವಿಷಯದಲ್ಲಿ ಹಾಗೆ ಮಾಡಲಾಗುತ್ತದೆ.

CuTeDSL ನಲ್ಲಿ, ಕೆಳಗಿನ ಉದಾಹರಣೆಯಿಂದ ವಿವರಿಸಿರುವಂತೆ, ಈ ನಿರ್ಧಾರಗಳನ್ನು ಹೋಸ್ಟ್-ಸೈಡ್ ಲಾಂಚ್ ಕೋಡ್‌ನಲ್ಲಿ ತೆಗೆದುಕೊಳ್ಳಬಹುದು ಮತ್ತು ಅಗತ್ಯವಿರುವಂತೆ ಗ್ರಿಡ್ ತಡೆಗೋಡೆಗಳೊಂದಿಗೆ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಸಂಯೋಜಿಸಬಹುದು ಮತ್ತು ಇಂಟರ್‌ಲೀವ್ ಮಾಡಬಹುದು. ಅಂತಹ ಹೈಬ್ರಿಡ್ ಕಾರ್ಯತಂತ್ರದೊಂದಿಗೆ, ಕನಿಷ್ಠ ಲೇಟೆನ್ಸಿಗಳನ್ನು ತಲುಪಲು ಡಿಕೋಡ್ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಆಪ್ಟಿಮೈಸ್ ಮಾಡಬಹುದು ಮತ್ತು ಗರಿಷ್ಠ ಥ್ರುಪುಟ್‌ಗಾಗಿ ಪ್ರಿಫಿಲ್ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಆಪ್ಟಿಮೈಸ್ ಮಾಡಬಹುದು. CUDA-ಆಧಾರಿತ ಪರಿಹಾರಗಳಿಗಿಂತ ಭಿನ್ನವಾಗಿ, ಈ ತರ್ಕವನ್ನು ಹೆಚ್ಚು ಸಂಕ್ಷಿಪ್ತವಾಗಿ ಮತ್ತು ಕಡಿಮೆ ಕಂಪೈಲ್-ಸಮಯದ ಓವರ್‌ಹೆಡ್‌ನೊಂದಿಗೆ ವ್ಯಾಖ್ಯಾನಿಸಬಹುದು.

python

@cute.kernel
def dispatch_kernel_decode(...):
	route_tokens(...)
	grid_barrier(grid_barrier_ptr)
	copy_tokens_to_experts(...)

@cute.kernel
def dispatch_kernel_prefill_route(...):
	route_tokens(...)
	
@cute.kernel
def dispatch_kernel_prefill_copy(...):
	copy_tokens_to_experts(...)
	
@cute.jit
def dispatch(...):
	num_route_blocks = ...
	num_copy_blocks = ...
	
	if num_copy_blocks < NUM_SMS:
		dispatch_kernel_decode(...).launch(
			grid=(min(max(num_route_blocks, num_copy_blocks), NUM_SMS), 1, 1)
		)
	else:
		dispach_kernel_prefill_route(...).launch(
			grid=(num_route_blocks, 1, 1)
		)
		dispatch_kernel_prefill_copy(...).launch(
			grid=(num_copy_blocks, 1, 1)
		)

ನಮ್ಮ MoE ಆಲ್-ಟು-ಆಲ್ ಕರ್ನೆಲ್‌ಗಳ ಸಂದರ್ಭದಲ್ಲಿ, ಕೆಳಗೆ ವಿವರವಾಗಿರುವಂತೆ, ಗ್ರಿಡ್ ತಡೆಗೋಡೆಗಳೊಂದಿಗೆ ಅಥವಾ ಇಲ್ಲದೆ ಕರ್ನೆಲ್‌ಗಳನ್ನು ನಿರ್ಮಿಸುವ ನಮ್ಯತೆಯನ್ನು ಹೊಂದಿರುವುದು ಮತ್ತು ಕಂಪೈಲ್ ಸಮಯದಲ್ಲಿ ಅವುಗಳನ್ನು ಆಪ್ಟಿಮೈಸ್ ಮಾಡುವುದು CuTeDSL ಮೂಲಕ ಗಣನೀಯ ಕಾರ್ಯಕ್ಷಮತೆ ಲಾಭವನ್ನು ಸಾಧಿಸಲು ನಮಗೆ ಸಹಾಯ ಮಾಡಿತು.

ನಾವು CuTeDSL ನಲ್ಲಿ ನಿರ್ಮಿಸಿದ ಮೊದಲ ಕರ್ನೆಲ್ ವಿಶೇಷ RMS ನಾರ್ಮ್ ಕರ್ನೆಲ್ ಆಗಿದೆ, ಇದು ಸಣ್ಣ ಡೈಮೆನ್ಶನ್‌ಗಳಿಗಾಗಿ ಆಪ್ಟಿಮೈಸ್ ಮಾಡಲಾಗಿದೆ. ಇತ್ತೀಚಿನ Qwen ಮಾದರಿಗಳ ಸೆಲ್ಫ್-ಅಟೆನ್ಷನ್ ಲೇಯರ್‌ನಲ್ಲಿ ಬಳಸಲಾಗುವ QK ನಾರ್ಮ್ ಅನ್ನು, ಪ್ರತಿ-ಹೆಡ್ ಆಧಾರದ ಮೇಲೆ, 64, 128 ಅಥವಾ 256 ಎಲಿಮೆಂಟ್‌ಗಳ ಮೇಲೆ ಅನ್ವಯಿಸಲಾಗುತ್ತದೆ. ಸಾವಿರಾರು ಮೀರುವ ಹಿಡನ್ ಡೈಮೆನ್ಶನ್‌ಗಳ ಮೇಲೆ ಕಾರ್ಯನಿರ್ವಹಿಸುವ ಸಾಮಾನ್ಯ RMS ನಾರ್ಮ್‌ಗಿಂತ ಭಿನ್ನವಾಗಿ, ಶೇರ್ಡ್ ಮೆಮೊರಿ ಮತ್ತು ರಿಜಿಸ್ಟರ್‌ಗಳು ಎರಡರ ಮೂಲಕ ಸಮತಾದ ಕಡಿತದ ಅಗತ್ಯವಿರುತ್ತದೆ, QK ನಾರ್ಮ್‌ಗೆ ವಾರ್ಪ್‌ಗಳಾದ್ಯಂತ ಮಾತ್ರ ಕಡಿತದ ಅಗತ್ಯವಿರುತ್ತದೆ, ಜೊತೆಗೆ ಬಹು ಹೆಡ್‌ಗಳನ್ನು ಒಂದೇ ವಾರ್ಪ್‌ಗೆ ಸ್ಕ್ವೀಜ್ ಮಾಡುವ ಸಾಧ್ಯತೆಯಿರುತ್ತದೆ. ಇದು ಎಲ್ಲಾ ಹೆಡ್ ಡೈಮೆನ್ಶನ್‌ಗಳು ಮತ್ತು ಡೇಟಾ ಪ್ರಕಾರಗಳಿಗೆ ಅತ್ಯುತ್ತಮ ಟೋಕನ್-ಟು-ವಾರ್ಪ್ ನಿಯೋಜನೆಗಳನ್ನು ಹುಡುಕಲು ನಾವು ಕಂಪೈಲ್ ಸಮಯದಲ್ಲಿ ಆಕ್ರಮಣಕಾರಿಯಾಗಿ ಆಪ್ಟಿಮೈಸ್ ಮಾಡಬಹುದಾದ ಬಳಕೆ ಪ್ರಕರಣವಾಯಿತು.

ವೆಕ್ಟರ್‌ನ ವರ್ಗೀಕೃತ ನಾರ್ಮ್ ಅನ್ನು ಲೆಕ್ಕಹಾಕಲು ಸಮತಾದ ಮೊತ್ತದ ಅಗತ್ಯವಿದ್ದರೂ, RMS ನಾರ್ಮ್ ತುಲನಾತ್ಮಕವಾಗಿ ನೇರವಾದ ಕಾರ್ಯಾಚರಣೆಯಾಗಿದೆ:

cpp

r = w * (x * ((x * x).sum(dim=-1) / x.shape[-1] + EPS))

ಅಟೆನ್ಷನ್ ಹೆಡ್‌ಗಳು ಸಾಮಾನ್ಯವಾಗಿ ಚಿಕ್ಕದಾಗಿರುತ್ತವೆ ಮತ್ತು ಪ್ರತಿ-ಥ್ರೆಡ್ ಕಾರ್ಯಾಚರಣೆಗಳು ಹೆಚ್ಚು ವೆಕ್ಟರೈಸ್ ಮಾಡಲ್ಪಟ್ಟಿರುವುದರಿಂದ, ಒಂದೇ ವಾರ್ಪ್ ವರ್ಗ ನಾರ್ಮ್ ಅನ್ನು ಲೆಕ್ಕಹಾಕಬಹುದು. ಮೆಮೊರಿ ಬ್ಯಾಂಡ್‌ವಿಡ್ತ್ ಅನ್ನು ಗರಿಷ್ಠಗೊಳಿಸಲು, ಪ್ರತಿ ಥ್ರೆಡ್ 128 ಬಿಟ್‌ಗಳ ಇನ್‌ಪುಟ್ ಅನ್ನು ಪ್ರಕ್ರಿಯೆಗೊಳಿಸಬೇಕು, ಇದು 8 bfloat16 ಮೌಲ್ಯಗಳಿಗೆ ಅನುವಾದಿಸುತ್ತದೆ. ಪರಿಣಾಮವಾಗಿ, ಒಂದು ಹೆಡ್ ಅನ್ನು ಪ್ರಕ್ರಿಯೆಗೊಳಿಸಲು ಹಾಫ್-ವಾರ್ಪ್ ಸಾಕು, ಮತ್ತು shuffle_sync_bfly ನಿಂದ ನಿರ್ವಹಿಸಲ್ಪಡುವ ಸಮತಾದ ಕಡಿತವನ್ನು ಬಳಸಿಕೊಂಡು ವರ್ಗ ನಾರ್ಮ್ ಅನ್ನು ನಿರ್ಧರಿಸಬಹುದು. ಹೆಡ್ ಡೈಮೆನ್ಶನ್ ಮತ್ತು ಡೇಟಾ ಪ್ರಕಾರದಿಂದ ಪ್ರಾರಂಭಿಸಿ, ಪ್ರತಿ ವಾರ್ಪ್‌ಗೆ ಹೆಡ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ನಾವು ನಿರ್ಧರಿಸಬಹುದು ಮತ್ತು ಕಂಪೈಲ್ ಮಾಡುವ ಸಮಯದಲ್ಲಿ ವಾದಗಳನ್ನು ಮತ್ತು ಕಡಿತ ಸುತ್ತುಗಳ ಸಂಖ್ಯೆಯನ್ನು ನಿರ್ದಿಷ್ಟಪಡಿಸಬಹುದು.

ಸ್ವಯಂಚಾಲಿತವಾಗಿ ವೆಕ್ಟರೈಸ್ ಮಾಡಿದ ಲೋಡ್‌ಗಳು/ಸ್ಟೋರ್‌ಗಳನ್ನು ಅವಲಂಬಿಸಿ ಮತ್ತು ವಾರ್ಪ್ ಕಡಿತಕ್ಕಾಗಿ ಕಡಿಮೆ-ಮಟ್ಟದ ಆರ್ಕಿಟೆಕ್ಚರ್-ನಿರ್ದಿಷ್ಟ ಕಾರ್ಯಾಚರಣೆಗಳನ್ನು ಬಳಸಿ, ನಾವು ನಂತರ ತುಲನಾತ್ಮಕವಾಗಿ ಕಡಿಮೆ ಸಾಲುಗಳ ಕೋಡ್‌ನೊಂದಿಗೆ ವಿಶೇಷ ಕರ್ನೆಲ್ ಅನ್ನು ನಿರ್ಮಿಸಬಹುದು. ಕರ್ನೆಲ್ ಯಾವ ಹೆಡ್ (Q ಅಥವಾ K) ಮೇಲೆ ಕೆಲಸ ಮಾಡಬೇಕು ಎಂಬುದನ್ನು ನಿರ್ಧರಿಸುತ್ತದೆ ಮತ್ತು ಅನ್ವಯಿಸಲು ಅನುಗುಣವಾದ ತೂಕಗಳನ್ನು ಗುರುತಿಸುತ್ತದೆ. ಇದು ಹೆಡ್‌ನಿಂದ ಸಂಪೂರ್ಣ ವೆಕ್ಟರ್ ಅನ್ನು ಲೋಡ್ ಮಾಡುತ್ತದೆ ಮತ್ತು ಮೊದಲು ಸ್ಥಳೀಯವಾಗಿ ಮೊತ್ತ ಮಾಡುತ್ತದೆ. ತರುವಾಯ, ಪ್ರತಿ ಹೆಡ್‌ಗೆ ಥ್ರೆಡ್‌ಗಳ ಸಂಖ್ಯೆಯ ಆಧಾರದ ಮೇಲೆ, ವರ್ಗ ನಾರ್ಮ್ ಅನ್ನು ಲೆಕ್ಕಹಾಕಲಾಗುತ್ತದೆ. ಅಂತಿಮವಾಗಿ, ಸ್ವಯಂಚಾಲಿತವಾಗಿ ವೆಕ್ಟರೈಸ್ ಮಾಡಲಾದ RMS ನಾರ್ಮ್ ಅನ್ನು ನಿರ್ಧರಿಸಲಾಗುತ್ತದೆ ಮತ್ತು ಮೆಮೊರಿಗೆ ಮತ್ತೆ ಬರೆಯಲಾಗುತ್ತದೆ. WEIGHT_BIAS ಮತ್ತು EPS ನಂತಹ ಮಾದರಿ-ನಿರ್ದಿಷ್ಟ ಪ್ಯಾರಾಮೀಟರ್‌ಗಳನ್ನು ಕಾನ್ಸ್ಟಂಟ್-ಫೋಲ್ಡ್ ಮಾಡಲಾಗುತ್ತದೆ.

python

@cute.kernel
def _qk_norm_kernel(
    q: cute.Tensor,
    k: cute.Tensor,
    q_weight: cute.Tensor,
    k_weight: cute.Tensor,
    WEIGHT_BIAS: cutlass.Constexpr[float],
    EPS: cutlass.Constexpr[float],
) -> None:
    # Find the block and thread indices.
    bidx, _, _ = cute.arch.block_idx()
    bdim, _, _ = cute.arch.block_dim()
    tidx, _, _ = cute.arch.thread_idx()

    _, num_q_heads, (threads_per_head, vecs_per_thread, vec_size) = q.shape
    _, num_k_heads, _ = k.shape
    head_dim = threads_per_head * vecs_per_thread * vec_size

    # Determine the thread and the vector to work on.
    vec_id = (bidx * bdim + tidx) // threads_per_head
    thread_id = tidx % threads_per_head

    # Determine the token and head indices.
    token = vec_id // (num_q_heads + num_k_heads)
    head_offset = vec_id % (num_q_heads + num_k_heads)
    is_q = head_offset < num_q_heads

    # Determine which tensor to work with.
    thread_coord = (thread_id, None, None)
    head_coord = (token, head_offset if is_q else head_offset - num_q_heads, thread_coord)

    # Load the head into registers.
    x_ref = q[head_coord] if is_q else k[head_coord]
    w = (q_weight if is_q else k_weight)[thread_coord].load()

    # Compute the squared sum locally.
    x = x_ref.load()
    x_f32 = x.to(cute.Float32)
    x_square = x_f32 * x_f32
    x_square_sum = x_square.reduce(cute.ReductionOp.ADD, init_val=0.0)

    # Sum up across the warp.
    mask = activemask()
    start = cute.arch.WARP_SIZE // (cute.arch.WARP_SIZE // threads_per_head * 2)
    for offset in [start >> i for i in range(5) if (start >> i) > 0]:
        x_square_sum += cute.arch.shuffle_sync_bfly(
            x_square_sum,
            offset=offset,
            mask=mask,
        )

    # RMS norm.
    rms = cute.math.rsqrt(x_square_sum / head_dim + EPS)
    x_norm = ((w + WEIGHT_BIAS) * (x * rms)).to(x.element_type)
    x_ref.store(x_norm)

ಸಾಮಾನ್ಯ RMS ನಾರ್ಮ್ ಕರ್ನೆಲ್‌ಗೆ ಹೋಲಿಸಿದರೆ, ಈ ವಿಶೇಷ ರೂಪಾಂತರವು ಪ್ರಿಫಿಲ್ ಮತ್ತು ಡಿಕೋಡ್ ಎರಡಕ್ಕೂ 2-3x ಉತ್ತಮ ಥ್ರುಪುಟ್ ಮತ್ತು ಲೇಟೆನ್ಸಿಯನ್ನು ನೀಡುತ್ತದೆ.

MoE GEMM ಕರ್ನೆಲ್‌ಗಳು ಸಂಬಂಧಿತ ಮ್ಯಾಟ್ರಿಕ್ಸ್ ಗುಣಾಕಾರಗಳನ್ನು ನಿರ್ವಹಿಸಿದ ನಂತರ ತೂಕದ ಸಂಗ್ರಹಣೆಯನ್ನು ನಿರ್ವಹಿಸುವ ಮೂಲಕ, MoE ಡಿಸ್ಪ್ಯಾಚ್ ಮತ್ತು ಕಂಬೈನ್ ಇನ್‌ಪುಟ್ ಟೋಕನ್‌ಗಳನ್ನು ಅವು ರೂಟ್ ಮಾಡಲಾದ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗಳಿಗೆ ಶಫಲ್ ಮಾಡುತ್ತವೆ. ವಿವಿಧ ಉತ್ಪನ್ನಗಳು ವಿವಿಧ ಗಾತ್ರದ MoE ಮಾದರಿಗಳನ್ನು ಅವಲಂಬಿಸಿರುವುದರಿಂದ, ಕಾರ್ಯಕ್ಷಮತೆಯು ನಿರ್ದಿಷ್ಟ ಆಸಕ್ತಿಯಾಗಿದೆ. ನಿಯೋಜನೆಯ ಪ್ರಕಾರಗಳ ಆಧಾರದ ಮೇಲೆ, ನಾವು ಆಪ್ಟಿಮೈಸ್ ಮಾಡಲು ವಿವಿಧ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಪ್ರತ್ಯೇಕಿಸುತ್ತೇವೆ:

TP=1: ಒಂದೇ GPU ಗಾಗಿ ಸ್ಥಳೀಯ ಕರ್ನೆಲ್. ಡಿಸ್ಪ್ಯಾಚ್ ಟೋಕನ್‌ಗಳನ್ನು ಸಂಬಂಧಿತ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗಳಿಗೆ ಶಫಲ್ ಮಾಡುತ್ತದೆ ಮತ್ತು ಕಂಬೈನ್ ಅವುಗಳನ್ನು ಸಂಗ್ರಹಿಸುತ್ತದೆ. H100 ನಲ್ಲಿ Qwen3 30B ನಿಯೋಜನೆಗಳಿಗಾಗಿ ಪ್ರಾಥಮಿಕವಾಗಿ ಬಳಸಲಾಗುತ್ತದೆ.

DP=1, TP ≠ 1: ಟೆನ್ಸರ್-ಪ್ಯಾರಲಲ್, ಸಿಂಗಲ್-ನೋಡ್ ಕರ್ನೆಲ್. ಇನ್‌ಪುಟ್ ಅನ್ನು ಪ್ರತಿಕೃತಿಗೊಳಿಸಲಾಗುತ್ತದೆ, ಡಿಸ್ಪ್ಯಾಚ್ ಪ್ರಸ್ತುತ ಶ್ರೇಣಿಗಾಗಿ ಟೋಕನ್‌ಗಳನ್ನು ಆಯ್ಕೆ ಮಾಡುತ್ತದೆ ಮತ್ತು ಕಂಬೈನ್ ಎಲ್ಲಾ ಶ್ರೇಣಿಗಳಾದ್ಯಂತ ಅವುಗಳನ್ನು ಸಂಗ್ರಹಿಸುತ್ತದೆ. NVLink ಅಥವಾ NVLS ಮೂಲಕ ಸಂಗ್ರಹಣೆ. ದೊಡ್ಡ ಪ್ರಿಫಿಲ್ಲರ್‌ಗಳಿಗಾಗಿ ಬಳಸಲಾಗುತ್ತದೆ.

DP ≠ 1, NVLink: ಐಚ್ಛಿಕ ಟೆನ್ಸರ್ ಪ್ಯಾರಲಲಿಸಂ ಹೊಂದಿರುವ ಡೇಟಾ-ಪ್ಯಾರಲಲ್ ಕರ್ನೆಲ್. ಪ್ರತಿ ಶ್ರೇಣಿಯು ನಿರ್ದಿಷ್ಟ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗಳನ್ನು ಹೋಸ್ಟ್ ಮಾಡುವ ಪ್ರತಿಯೊಂದು ಇತರ ಶ್ರೇಳಿಗೆ ಟೋಕನ್‌ಗಳನ್ನು ಕಳುಹಿಸುತ್ತದೆ. ರೂಟಿಂಗ್ ಮಾಹಿತಿಯನ್ನು ಒಟ್ಟುಗೂಡಿಸಲಾಗುತ್ತದೆ, NVIDIA NVLink ಮೂಲಕ ಟೋಕನ್‌ಗಳನ್ನು ವರ್ಗಾಯಿಸಲಾಗುತ್ತದೆ. ಡೇಟಾ-ಪ್ಯಾರಲಲ್ H200 ಅಥವಾ GB200 ನಿಯೋಜನೆಗಳಿಗಾಗಿ ಬಳಸಲಾಗುತ್ತದೆ.

DP ≠ 1, InfiniBand: NVLink ಮತ್ತು NVIDIA ಕ್ವಾಂಟಮ್ InfiniBand ಎರಡರ ಮೂಲಕ ಕಾರ್ಯನಿರ್ವಹಿಸುವ ಪೀರ್-ಟು-ಪೀರ್ ಡೇಟಾ-ಪ್ಯಾರಲಲ್ ಮತ್ತು ಟೆನ್ಸರ್-ಪ್ಯಾರಲಲ್ ಕರ್ನೆಲ್. ಬಹು-ನೋಡ್ Hopper ನಿಯೋಜನೆಗಳಿಗಾಗಿ ಬಳಸಲಾಗುತ್ತದೆ.

ಸ್ಥಳೀಯ ಕರ್ನೆಲ್ ರೂಟಿಂಗ್ ಕೋಷ್ಟಕದ ಆಧಾರದ ಮೇಲೆ ಪ್ರತಿ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗೆ ನಿಯೋಜಿಸಲಾದ ಟೋಕನ್‌ಗಳ ಸಂಖ್ಯೆಯನ್ನು ಎಣಿಕೆ ಮಾಡುತ್ತದೆ ಮತ್ತು ಪ್ರತಿ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗೆ ಸೇರಿದ ಟೋಕನ್‌ಗಳ ಪಟ್ಟಿಯೊಳಗೆ ಪ್ರತಿ ಟೋಕನ್‌ಗೆ ಅನನ್ಯ ಸೂಚ್ಯಂಕವನ್ನು ನಿಯೋಜಿಸುತ್ತದೆ. ರೂಟಿಂಗ್ ಹಂತದ ನಂತರ, MoE GEMM ಕರ್ನೆಲ್‌ಗೆ ರವಾನಿಸಲು ಟೋಕನ್‌ಗಳನ್ನು ಐಚ್ಛಿಕವಾಗಿ ಪ್ಯಾಡ್ ಮಾಡಲಾದ ಬಫರ್‌ಗಳಿಗೆ ನಕಲು ಮಾಡಲಾಗುತ್ತದೆ. ಟೋಕನ್ ಸ್ಥಳಗಳನ್ನು ಬಳಸಿಕೊಂಡು, ಕಂಬೈನ್ ಕರ್ನೆಲ್ ನಂತರ ರಿಜಿಸ್ಟರ್‌ಗಳಲ್ಲಿ ಸಂಗ್ರಹಣೆಯನ್ನು ನಿರ್ವಹಿಸುತ್ತದೆ, ಅಂತಿಮ ಟೋಕನ್‌ಗಳನ್ನು ಔಟ್‌ಪುಟ್ ಬಫರ್‌ಗೆ ಬರೆಯುತ್ತದೆ.

ಪ್ರಿಫಿಲ್ ಮತ್ತು ಡಿಕೋಡ್ ನಡುವೆ ಡಿಸ್ಪ್ಯಾಚ್ ಹಂತವು ಬದಲಾಗುತ್ತದೆ: ಕೆಲವು ಟೋಕನ್‌ಗಳಿಗೆ, ಓದಲು ಸ್ವಲ್ಪ ರೂಟಿಂಗ್ ಮಾಹಿತಿ ಇರುತ್ತದೆ, ಆದ್ದರಿಂದ ನಾವು ಅದನ್ನು ಒಂದೇ ಬ್ಲಾಕ್‌ನಲ್ಲಿ ಪ್ರಕ್ರಿಯೆಗೊಳಿಸುತ್ತೇವೆ ಮತ್ತು ಗ್ರಿಡ್‌ನ ಉಳಿದ ಭಾಗಕ್ಕೆ ಸಂಕೇತಿಸುತ್ತೇವೆ, ವೈಯಕ್ತಿಕ ಬ್ಲಾಕ್‌ಗಳು ಟೋಕನ್‌ಗಳ ಚಂಕ್‌ಗಳನ್ನು ನಕಲು ಮಾಡುತ್ತವೆ. ಶೇರ್ಡ್ ಮೆಮೊರಿಯಲ್ಲಿ ಕೌಂಟರ್‌ಗಳನ್ನು ಪರಮಾಣು ರೀತಿಯಲ್ಲಿ ಹೆಚ್ಚಿಸುವ ಮೂಲಕ ಟೋಕನ್ ನಿಯೋಜನೆಯನ್ನು ನಿರ್ಧರಿಸಲಾಗುತ್ತದೆ. ಪ್ರಿಫಿಲ್ ಹೆಚ್ಚಿನ ರೂಟಿಂಗ್ ಮಾಹಿತಿಯನ್ನು ಪ್ರಕ್ರಿಯೆಗೊಳಿಸಬೇಕುವಾದ್ದರಿಂದ, ಪ್ರತ್ಯೇಕ ಕರ್ನೆಲ್‌ನಲ್ಲಿ ಬಹು ಬ್ಲಾಕ್‌ಗಳಾದ್ಯಂತ ರೂಟಿಂಗ್ ಅನ್ನು ವಿಂಗಡಿಸಲಾಗಿದೆ. ಶೇರ್ಡ್ ಮೆಮೊರಿಯಲ್ಲಿ ಕೌಂಟರ್‌ಗಳನ್ನು ಬಳಸಿ ಸೂಚ್ಯಂಕಗಳನ್ನು ಮೊದಲು ಸ್ಥಳೀಯವಾಗಿ ನಿಯೋಜಿಸಲಾಗುತ್ತದೆ ಮತ್ತು ನಂತರ ಪ್ರತಿ ಟೋಕನ್‌ನ ಸ್ಥಾನವನ್ನು ನಿರ್ಧರಿಸಲು ಗ್ಲೋಬಲ್ ಮೆಮೊರಿಯಲ್ಲಿ ಎಲ್ಲಾ ಬ್ಲಾಕ್‌ಗಳಾದ್ಯಂತ ಒಟ್ಟುಗೂಡಿಸಲಾಗುತ್ತದೆ. ಗ್ರಿಡ್-ವೈಸ್ ತಡೆಗೋಡೆಗಳನ್ನು ಬಳಸಲು ಸಾಧ್ಯವಾಗುವಂತೆ ಈ ಕರ್ನೆಲ್‌ನ ಗ್ರಿಡ್ ಅನ್ನು SM ಗಳ ಸಂಖ್ಯೆಗೆ ಕ್ಲ್ಯಾಂಪ್ ಮಾಡಲಾಗುತ್ತದೆ. ಲಭ್ಯವಿರುವ ಮೆಮೊರಿ ಬ್ಯಾಂಡ್‌ವಿಡ್ತ್ ಅನ್ನು ಸಂಪೂರ್ಣವಾಗಿ ಬಳಸಿಕೊಳ್ಳಲು ಹೆಚ್ಚು ದೊಡ್ಡ ಸಂಖ್ಯೆಯ ಬ್ಲಾಕ್‌ಗಳು ಮತ್ತು ವಾರ್ಪ್‌ಗಳಾದ್ಯಂತ ಲಾಂಚ್ ಮಾಡಲಾದ ನಂತರದ ಕರ್ನೆಲ್ ನಕಲುಗಳನ್ನು ನಿರ್ವಹಿಸುತ್ತದೆ.

combine kernel ತುಲನಾತ್ಮಕವಾಗಿ ಸರಳವಾಗಿದೆ, ಏಕೆಂದರೆ ಅದು vectorized accumulation ಅನ್ನು ನಿರ್ವಹಿಸುತ್ತದೆ, ಮತ್ತು routing information ಹಾಗೂ dispatch ಮೂಲಕ ಗಣನೆ ಮಾಡಲಾದ offsets ಮೇಲೆ ಅವಲಂಬಿತವಾಗಿರುತ್ತದೆ. ಇಲ್ಲಿ CuTeDSL ಬಹಳ ಉಪಯುಕ್ತವೆಂದು ಸಾಬೀತಾಯಿತು, ಏಕೆಂದರೆ ಗಣನೆಯನ್ನು ಸರಳ CuTe tensorsಗಳ ಮೂಲಕ ವ್ಯಕ್ತಪಡಿಸಬಹುದಾಗಿತ್ತು; ಅವು ಎಲ್ಲಾ data types ಮತ್ತು hidden dimensionsಗಳಿಗೆ ಪರಿಣಾಮಕಾರಿ code ಆಗಿ compile ಆಗುತ್ತವೆ.

ಟೆನ್ಸರ್-ಪ್ಯಾರಲಲ್ ನಿಯೋಜನೆಯಲ್ಲಿ, ರೂಟಿಂಗ್ ಮಾಹಿತಿ ಮತ್ತು ಇನ್‌ಪುಟ್ ಟೋಕನ್‌ಗಳನ್ನು ಪ್ರತಿಕೃತಿಗೊಳಿಸಲಾಗುತ್ತದೆ. ಪ್ರಸ್ತುತ ಶ್ರೇಣಿಯಲ್ಲಿರುವ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗಳಿಗೆ ನಿಯೋಜಿಸಲಾದ ಟೋಕನ್‌ಗಳನ್ನು ಮಾತ್ರ ರೂಟ್ ಮಾಡಲು ಮತ್ತು ಆಯ್ಕೆ ಮಾಡಲು ಸ್ವಲ್ಪ ಬದಲಾವಣೆಯೊಂದಿಗೆ ಸ್ಥಳೀಯ ಡಿಸ್ಪ್ಯಾಚ್ ಕರ್ನೆಲ್ ಅನ್ನು ಸಂಪೂರ್ಣವಾಗಿ ಮರುಬಳಕೆ ಮಾಡಬಹುದು. ಕೆಲವು ಪ್ಯಾರಾಮೀಟರ್‌ಗಳು ಮತ್ತು ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಪೆಷಲೈಸೇಶನ್‌ನೊಂದಿಗೆ, ಸ್ಥಳೀಯ ಡಿಸ್ಪ್ಯಾಚ್ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಮರುಬಳಕೆ ಮಾಡಲಾಗುತ್ತದೆ.

ಶ್ರೇಣಿಗಳ ಸಂಖ್ಯೆಯು ಸಾಮಾನ್ಯವಾಗಿ ಟೋಕನ್ ರೂಟ್ ಮಾಡಲಾದ ಎಕ್ಸ್‌ಪರ್ಟ್‌ಗಳ ಸಂಖ್ಯೆಗಿಂತ ಕಡಿಮೆಯಿರುವುದರಿಂದ, ಸರಾಸರಿ ಸಂದರ್ಭದಲ್ಲಿ ಪ್ರತಿ ಟೋಕನ್ ಪ್ರತಿ ಶ್ರೇಳಿಗೆ ರೂಟ್ ಮಾಡಲಾದ ತನ್ನ ಪ್ರತಿಕೃತಿಗಳಲ್ಲಿ ಒಂದನ್ನು ಹೊಂದಿದೆ ಎಂದು ಊಹಿಸಬಹುದು. ಈ ಪರಿಸ್ಥಿತಿಯಲ್ಲಿ, ಟೋಕನ್‌ನ ಸ್ಥಳೀಯ ಪ್ರತಿಗಳನ್ನು ಒಟ್ಟುಗೂಡಿಸಲು ಸ್ಥಳೀಯ ಕಂಬೈನ್ ಕರ್ನೆಲ್ ಅನ್ನು ಮರುಬಳಕೆ ಮಾಡಬಹುದು, ನಂತರ ಭಾಗವಹಿಸುವ ಶ್ರೇಣಿಗಳಾದ್ಯಂತ ಮೊತ್ತ ಮಾಡಲು ಆಲ್-ರೆಡ್ಯೂಸ್ ಮಾಡಬಹುದು. ಪ್ರಿಫಿಲ್ ಕರ್ನೆಲ್‌ನ ಸಂದರ್ಭದಲ್ಲಿ, ನಾವು ಪ್ರತ್ಯೇಕ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಲಾಂಚ್ ಮಾಡುತ್ತೇವೆ, ಆದಾಗ್ಯೂ ಡಿಕೋಡ್ ಎರಡೂ ಕಾರ್ಯಾಚರಣೆಗಳನ್ನು ಒಂದೇ ಲಾಂಚ್‌ನಲ್ಲಿ ಫ್ಯೂಸ್ ಮಾಡುತ್ತದೆ.

ಆಲ್-ರೆಡ್ಯೂಸ್ಗಾಗಿ, ಕೇವಲ ಎರಡು ಶ್ರೇಣಿಗಳು ಭಾಗವಹಿಸಿದರೆ ನಾವು ಸರಳ NVLink ಅನ್ನು ಬಳಸುತ್ತೇವೆ, ಥ್ರುಪುಟ್‌ನಲ್ಲಿ 10-20% ಹೆಚ್ಚಳಕ್ಕಾಗಿ 4 ಅಥವಾ ಹೆಚ್ಚಿನ ಶ್ರೇಣಿಗಳು involucrated ಇರುವಾಗ NVLS ಗೆ ಬದಲಾಯಿಸುತ್ತೇವೆ. DSL ಗೆ ಧನ್ಯವಾದಗಳು, MoE ಮತ್ತು ಸಾಮಾನ್ಯ ಆಲ್-ರೆಡ್ಯೂಸ್ ಕರ್ನೆಲ್‌ಗಳ ನಡುವೆ ಹೆಚ್ಚಿನ ಕಡಿತ ತರ್ಕ ಮತ್ತು ಪೀರ್ ಬಫರ್ ವಿನಿಮಯವನ್ನು ನಾವು ಹಂಚಿಕೊಳ್ಳಬಹುದು. NVLS ಗಾಗಿ multimem.ld_reduce ಮತ್ತು multimem.st ಮೇಲಿನ DSL ಕಾರ್ಯಾಚರಣೆಗಳು ಅಗತ್ಯವಿರುವ ಕಾರಣ, ಹಾರ್ಡ್‌ವೇರ್ ಮೇಲೆ ನಿಕಟ ನಿಯಂತ್ರಣವು ಪ್ರಯೋಜನಕಾರಿಯಾಗಿದೆ. ಕಡಿತದ ಅನುಷ್ಠಾನವು ಹೆಚ್ಚಾಗಿ NCCL ಪರೀಕ್ಷೆಗಳಿಂದ all_reduce ಉದಾಹರಣೆಯನ್ನು ಅನುಸರಿಸುತ್ತದೆ.

ಡೇಟಾ ಪ್ಯಾರಲಲ್ ಕರ್ನೆಲ್‌ಗಳು ಬಹು ಸಾಧನಗಳಿಗೆ ಟೋಕನ್‌ಗಳನ್ನು ಒಟ್ಟುಗೂಡಿಸಲು NVLink ಅನ್ನು ಅವಲಂಬಿಸಿವೆ. ಡಿಸ್ಪ್ಯಾಚ್ ಹಂತದಲ್ಲಿ, ಟೋಕನ್‌ಗಳನ್ನು ವಿನಿಮಯ ಮಾಡಿಕೊಳ್ಳುವ ಮೊದಲು, ಪೀರ್‌ಗಳು ತಾವು ಕಳುಹಿಸುತ್ತಿರುವ ಪ್ರತಿ ಟೋಕನ್‌ಗೆ ಪ್ರತಿ ಶ್ರೇಣಿಯಲ್ಲಿ ಗುರಿ ಆಫ್‌ಸೆಟ್ ಅನ್ನು ಗುರುತಿಸಲು ಮೊದಲು ಸ್ಥಳೀಯವಾಗಿ ರೂಟ್ ಮಾಡುತ್ತವೆ. ತರುವಾಯ, ಸಂಘರ್ಷಗಳನ್ನು ತಪ್ಪಿಸಲು ಪ್ರತಿ ಕಳುಹಿಸುವವರಿಗೆ ಹಂಚಿಕೆ ಮಾಡಲಾದ ಸ್ವೀಕರಿಸುವ ಬಫರ್‌ಗಳಿಗೆ ಅವರು ಟೋಕನ್‌ಗಳನ್ನು ನಕಲು ಮಾಡಲು ಪ್ರಾರಂಭಿಸುತ್ತಾರೆ. ಪೂರ್ಣಗೊಳ್ಳುವುದನ್ನು ಸಂಕೇತಿಸಲು ಪ್ರತಿ-ಟೋಕನ್ ರಿಲೀಸ್-ಸ್ಟೋರ್ ಬರವಣಿಗೆಗಳನ್ನು ಬಳಸಲಾಗುತ್ತದೆ. ಸ್ವೀಕರಿಸುವವರ ಕೊನೆಯಲ್ಲಿ, MoE GEMM ಗಾಗಿ ಎಲ್ಲಾ ಟೋಕನ್‌ಗಳನ್ನು ಒಂದೇ ಬಫರ್‌ಗೆ ಸಂಕುಚಿಸಲು ಸಾಧ್ಯವಾಗುವಂತೆ ರೂಟಿಂಗ್ ಮಾಹಿತಿಯನ್ನು ಒಟ್ಟುಗೂಡಿಸಲಾಗುತ್ತದೆ. ಬರವಣಿಗೆಗಳು ಪೂರ್ಣಗೊಳ್ಳಲು ಕಾಯಲು ಸ್ವೀಕರಿಸುವವರು ಅಕ್ವೈರ್-ಲೋಡ್‌ನೊಂದಿಗೆ ಪೀರ್-ಬರೆದ ಫ್ಲಾಗ್ ಅನ್ನು ಪೋಲ್ ಮಾಡುತ್ತಾರೆ.

combine ಬದಿಯಲ್ಲಿ, ಟೋಕನ್‌ಗಳನ್ನು NVLink ಮೂಲಕ source ranksಗಳಿಗೆ ನಕಲಿಸಲಾಗುತ್ತದೆ, ಮತ್ತು ಪ್ರತಿ ಟೋಕನ್‌ಗೆ ಸ್ವೀಕೃತಿಯನ್ನು ಸೂಚಿಸುವ flag ಇರುತ್ತದೆ. combine kernelವು local ಮತ್ತು singlenode ರೂಪಾಂತರಗಳಿಗೆ ಸಮಾನವಾಗಿದ್ದು, addressing‌ಗೆ ಸಂಬಂಧಿಸಿದ ಕೆಲವು ಸಣ್ಣ ವ್ಯತ್ಯಾಸಗಳು ಮತ್ತು ಹೆಚ್ಚುವರಿ synchronization logic ಅನ್ನು ಹೊಂದಿದೆ.

ಈ ಕರ್ನೆಲ್ ಅನ್ನು ಕಾರ್ಯಗತಗೊಳಿಸಲು ಸಿಸ್ಟಮ್ ಮತ್ತು ಸಾಧನ ಮಟ್ಟದಲ್ಲಿ ಸ್ಕೋಪ್ ಮಾಡಲಾದ ರಿಲೀಸ್-ಅಕ್ವೈರ್ ಲೋಡ್‌ಗಳು ಮತ್ತು ಸ್ಟೋರ್‌ಗಳನ್ನು ಬಹಿರಂಗಪಡಿಸಲು ನಾವು PTX ಸುತ್ತುವರೆದಿರುವವರ ಲೈಬ್ರರಿಯನ್ನು ನಿರ್ಮಿಸಿದ್ದೇವೆ. ಸಿಂಕ್ರೊನೈಸೇಶನ್ ಫ್ಲಾಗ್‌ಗಳ ಗ್ರಿನ್ಯುಲಾರಿಟಿಯನ್ನು ಸುಧಾರಿಸುವ ಮೂಲಕ ಮತ್ತು NVLink ಮೂಲಕ ಟೋಕನ್‌ಗಳನ್ನು ಕಳುಹಿಸಲು ಮತ್ತು ಸ್ವೀಕರಿಸಲು ಹೆಚ್ಚಿನ ಬ್ಲಾಕ್‌ಗಳನ್ನು ಹಂಚಿಕೆ ಮಾಡುವ ಮೂಲಕ, ನಮ್ಮ ಹಿಂದಿನ CUDA ಕರ್ನೆಲ್‌ಗಳಿಗೆ ಹೋಲಿಸಿದರೆ ನಾವು 2x ಸುಧಾರಣೆಗಳನ್ನು ಸಾಧಿಸಿದ್ದೇವೆ.

ನಮ್ಮ Peer-to-Peer kernelsಗಳು ಹಿಂದೆ Hopper GPUs ಮತ್ತು ConnectX-7 InfiniBand adapters ಜೊತೆಗೆ state-of-the-art ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಸಾಧಿಸಿದ್ದವು. ನಂತರ, ಹೆಚ್ಚಿನ ಕಾರ್ಯಕ್ಷಮತೆ ಲಾಭಗಳಿಗಾಗಿ ನಾವು combine kernelsಗಳನ್ನು CuTeDSLಗೆ port ಮಾಡಿದೆವು. CUDA ಆವೃತ್ತಿಗಳು ಈಗಾಗಲೇ compile-time specialization ಮೇಲೆ ಬಹಳ ಅವಲಂಬಿತವಾಗಿದ್ದುದರಿಂದ, ಅವು ನಮ್ಮ repositoryಯಲ್ಲಿ compilation-time bottlenecks ಆಗಿ ಪರಿಣಮಿಸಿದವು ಮತ್ತು ಕೆಲವು parameters ಅನ್ನು ಕೈಬಿಡುವಂತೆ ಮಾಡಿತು. JIT-compiled DSLಗೆ migrate ಮಾಡುವ ಮೂಲಕ, ನಾವು ಹೆಚ್ಚು parametersಗಾಗಿ specialize ಮಾಡಿದೆವು ಮತ್ತು grid allocations ಅನ್ನು fine-tune ಮಾಡಿದೆವು, latencyಗೆ ಸುಮಾರು 3% boost ಪಡೆಯುವ ಸಲುವಾಗಿ prefill ಮತ್ತು decode ನಡುವೆ ಇನ್ನಷ್ಟು ವ್ಯತ್ಯಾಸಗಳನ್ನು ಪರಿಚಯಿಸಿದೆವು.

ROSE ಕರ್ನೆಲ್‌ಗಳ ದೃಷ್ಟಿಕೋನದಿಂದ, Blackwell ಬೆಂಬಲವು ಎರಡು ವರ್ಗಗಳಾಗಿ ವಿಂಗಡಿಸಲ್ಪಟ್ಟಿದೆ: ಆರ್ಕಿಟೆಕ್ಚರ್-ಅಜ್ಞೋಸ್ಟಿಕ್ ಮತ್ತು ಆರ್ಕಿಟೆಕ್ಚರ್-ನಿರ್ದಿಷ್ಟ ಕರ್ನೆಲ್‌ಗಳು. ಟೆನ್ಸರ್ ಕೋರ್‌ಗಳನ್ನು ಅವಲಂಬಿಸಿರುವ ಆರ್ಕಿಟೆಕ್ಚರ್-ನಿರ್ದಿಷ್ಟ ಕರ್ನೆಲ್‌ಗಳಿಗೆ ಹೊಂದಿಕೊಳ್ಳಲು ಸ್ಪಷ್ಟ ಕೆಲಸದ ಅಗತ್ಯವಿರುತ್ತದೆ. ಆದಾಗ್ಯೂ, DSL ನ ಬಳನೆಗೆ ಧನ್ಯವಾದಗಳು, ಆರ್ಕಿಟೆಕ್ಚರ್-ಅಜ್ಞೋಸ್ಟಿಕ್ ಕರ್ನೆಲ್‌ಗಳು ಸುಲಭವಾಗಿ ಹೊಂದಿಕೊಳ್ಳುತ್ತವೆ.

Hopper ನ 132 ಕ್ಕೆ ವ್ಯತಿರಿಕ್ತವಾಗಿ Blackwell GPU ವರೆಗೆ 160 SM ಗಳನ್ನು ಹೊಂದಿದೆ. ಹೆಚ್ಚಿನ ಕರ್ನೆಲ್‌ಗಳಿಗೆ, ಇದು ಕಂಪೈಲ್-ಸಮಯದ ಸ್ಥಿರಾಂಕವಾಗಿದೆ ಮತ್ತು ರನ್ ಮಾಡಲು ಅತ್ಯುತ್ತಮ (ಪ್ರಿಫಿಲ್ ವಿರುದ್ಧ ಡಿಕೋಡ್) ಕರ್ನೆಲ್ ಅನ್ನು ನಿರ್ಧರಿಸಲು ಹೋಸ್ಟ್-ಸೈಡ್ ಲಾಂಚರ್ ಕೋಡ್‌ನಲ್ಲಿರುವ ತರ್ಕವು ಇದನ್ನು ಗಣನೆಗೆ ತೆಗೆದುಕೊಳ್ಳುತ್ತದೆ. ಹೆಚ್ಚುವರಿಯಾಗಿ, Blackwell 256-ಬಿಟ್ ಲೋಡ್ ಸೂಚನೆಗಳನ್ನು ಪರಿಚಯಿಸುತ್ತದೆ, ಮೆಮೊರಿ ಬ್ಯಾಂಡ್‌ವಿಡ್ತ್ ಬಳಕೆಯನ್ನು ಸುಧಾರಿಸಲು ವೆಕ್ಟರ್ ಅಗಲವನ್ನು (CuTe ಟೆನ್ಸರ್‌ಗಳ ಕೊನೆಯ ಡೈಮೆನ್ಶನ್) ಹೆಚ್ಚಿಸುವ ಮೂಲಕ ಇದನ್ನು ಸುಲಭವಾಗಿ ಹೊರಸೂಸಬಹುದು.

Blackwell ಗಮನಾರ್ಹ ಸಂಖ್ಯೆಯ ಕ್ವಾಂಟೈಸೇಶನ್ ಯೋಜನೆಗಳಿಗೆ (MXFP8, MXFP4, NVFP4) ಹಾರ್ಡ್‌ವೇರ್ ಬೆಂಬಲವನ್ನು ಪರಿಚಯಿಸುತ್ತದೆ. FP8 ಬ್ಲಾಕ್-ವೈಸ್ ಸ್ಕೇಲಿಂಗ್‌ಗಾಗಿ ಅಸ್ತಿತ್ವದಲ್ಲಿರುವ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಟ್ವೀಕ್ ಮಾಡುವ ಮೂಲಕ ಇವುಗಳ ಬೆಂಬಲವನ್ನು ಪ್ರಾಥಮಿಕವಾಗಿ ಕಾರ್ಯಗತಗೊಳಿಸಲಾಯಿತು. ಕಂಪೈಲ್-ಸಮಯದ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್ ಅನ್ನು ಸ್ಫೋಟಿಸದೆ ಹೊಸ ಮೋಡ್‌ಗಳನ್ನು ಬೆಂಬಲಿಸಬಹುದಾದ್ದರಿಂದ JIT ಬೆಂಬಲವು ಇಲ್ಲಿ ನಿರ್ಣಾಯಕವಾಗಿದೆ. ಹಿಂದೆ, ಕೆಲವು ಎಲಿಮೆಂಟ್‌ಗಳಾದ್ಯಂತ, ವಾರ್ಪ್‌ನ ಅರ್ಧ ಅಥವಾ ಕಾಲು ಭಾಗದಷ್ಟು ಸಮತಾದ ಕಡಿತಗಳನ್ನು ವ್ಯಕ್ತಪಡಿಸಲು Triton ಅಹಿತಕರವಾಗಿದೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ. NVFP4 ಮತ್ತು MXFP8/MXFP4 ಗಾಗಿ ಕ್ರಮವಾಗಿ 16 ಅಥವಾ 32 ಎಲಿಮೆಂಟ್‌ಗಳಾದ್ಯಂತ ಸಂಪೂರ್ಣ ಗರಿಷ್ಠ ಮೌಲ್ಯವನ್ನು ಲೆಕ್ಕಹಾಕುವಾಗ ಅತ್ಯುತ್ತಮ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಸಾಧಿಸಲು ಕಡಿತಗಳ ಮೇಲೆ ಉತ್ತಮ-ಗ್ರಿನ್ಯುಲಾರಿಟಿ ನಿಯಂತ್ರಣವು ನಿರ್ಣಾಯಕವಾಗಿತ್ತು.

ತುಲನಾತ್ಮಕವಾಗಿ ಹೊಸ ಭಾಷೆಯಾಗಿರುವುದರಿಂದ, LLM ಗಳು ತಮ್ಮ ಪ್ರೀ-ಟ್ರೈನಿಂಗ್ ಸಮಯದಲ್ಲಿ CuTeDSL ಕೋಡ್‌ಗೆ ಕಡಿಮೆ ಮಾನ್ಯತೆಯನ್ನು ಹೊಂದಿದ್ದವು. ಹೆಚ್ಚುವರಿಯಾಗಿ, TMA ಮತ್ತು ಆರ್ಕಿಟೆಕ್ಚರ್-ನಿರ್ದಿಷ್ಟ ಮ್ಯಾಟ್ರಿಕ್ಸ್ ಗುಣಾಕಾರ ಸೂಚನೆಗಳಿಗಾಗಿ ಹೆಚ್ಚು ಆಪ್ಟಿಮೈಸ್ ಮಾಡಲಾದ ಅಸ್ತಿತ್ವದಲ್ಲಿರುವ CUTLASS ಅಥವಾ Quack ಉದಾಹರಣೆಗಳಿಂದ, ಆಕ್ಟಿವೇಷನ್‌ಗಳು ಮತ್ತು ನಾರ್ಮ್‌ಗಳಿಗಾಗಿ ಆಪ್ಟಿಮೈಸ್ ಮಾಡಲಾದ ಕರ್ನೆಲ್‌ಗಳನ್ನು ಹೊಂದಿರುವ ನಮ್ಮ ಪ್ರಸ್ತುತ ಕೋಡ್‌ನ ಶೈಲಿಯು ಸ್ವಲ್ಪ ಭಿನ್ನವಾಗಿತ್ತು. LLM ಗಳು ಕರ್ನೆಲ್ ಬರೆಯುವಲ್ಲಿ ಸಹಾಯ ಮಾಡುವ ಹಂತವನ್ನು ತಲುಪುವುದು ಕ್ರಮೇಣ ಪ್ರಕ್ರಿಯೆಯಾಗಿತ್ತು.

ಕರ್ಬಣ ಪರಿವರ್ತನೆಗಾಗಿ agentsಗಳ ಮೇಲೆ ಅವಲಂಬಿಸಬಹುದಾದ ಮೊದಲು, ನಾವು ಸುಸಂಗತ patternsಗಳನ್ನು ಅನುಸರಿಸಿ kernels ಮತ್ತು utility functionsಗಳ ಆರಂಭಿಕ corpus ಅನ್ನು ಕೈಯಾರೆ ನಿರ್ಮಿಸಿಕೊಂಡೆವು. corpus ಮುಂದುವರಿದಂತೆ, ನಾವು CUDAಯಿಂದ CuTeDSLಗೆ rough translationsಗಳನ್ನು ಮಾಡಬಹುದಾಗಿತ್ತು, ನಂತರ ಗಮನಾರ್ಹ manual rework ಅಗತ್ಯವಾಗುತ್ತಿತ್ತು. ವಿಶೇಷವಾಗಿ, pointer-to-tensor conversions ಮತ್ತು vectorizationನಲ್ಲಿ modelsಗೆ ತೊಂದರೆ ಆಗುವುದನ್ನು ನಾವು ಕಂಡೆವು; ಪರೀಕ್ಷೆಯನ್ನು compile ಮಾಡಲು ಯತ್ನಿಸುವಾಗ type errors ಸರಿಪಡಿಸಲು ಅವು ಅಟ್ಟಿಹೋಗುತ್ತಿದ್ದವು. kernelsಗಳ grid-based ಸ್ವಭಾವದ ಬಗ್ಗೆ ಅವುಗಳಲ್ಲಿ ಅಲ್ಪ ಅರ್ಥಮಾಡಿಕೊಳ್ಳುವಿಕೆಯಷ್ಟೇ ಕಂಡುಬಂದಿತು, thread ಮತ್ತು grid synchronization ಅನ್ನು ಗೊಂದಲಗೊಳಿಸುತ್ತಿದ್ದವು. ಅತ್ಯಂತ ನಿರಾಶಾದಾಯಕವಾಗಿ, ನಮ್ಮದೇ CuTeDSL replacementsಗಳನ್ನು ನಿರ್ಮಿಸುವ ಬದಲು CUDA-specific librariesಗಳ ಅಸ್ತಿತ್ವದಲ್ಲೇ ಇಲ್ಲದ translationsಗಳನ್ನು hallucinate ಮಾಡುವ ಪ್ರಲೋಭನೆ ಇತ್ತು. ಕಾಲಕ್ರಮೇಣ, CuTeDSL kernelsಗಳ corpus ವೃದ್ಧಿಯಾಗುತ್ತಿದ್ದಂತೆ, LLMಗಳು ಬಹುಪಾಲು skeleton code ಅನ್ನು ನಿರ್ಮಿಸಿ ಖಾಲಿ ಜಾಗಗಳನ್ನು ತುಂಬಬಹುದೆಂದು ನಾವು ಕಂಡುಹಿಡಿದೆವು, ಇದರಿಂದ high-level design ಮೇಲೆ ನಾವು ಗಮನಹರಿಸಬಹುದು.

CUDA ಗೆ ಹೋಲಿಸಿದರೆ, CuTeDSL ಹೆಚ್ಚು ಉತ್ತಮ ಕಂಪೈಲ್-ಸಮಯದ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ತಲುಪಿಸುತ್ತದೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ, ಇದು ಹೆಚ್ಚಿನ ಪ್ಯಾರಾಮೀಟರ್‌ಗಳಿಗಾಗಿ ಸ್ಪೆಷಲೈಸ್ ಮಾಡಲು ಮತ್ತು ಕಾರ್ಯಕ್ಷಮತೆ ಸುಧಾರಣೆಗಳನ್ನು ಸಾಧಿಸಲು ನಮಗೆ ಅನುವು ಮಾಡಿಕೊಡುತ್ತದೆ. Triton ಗೆ ಹೋಲಿಸಿದರೆ ಕೆಲವು ಅಮಸ್ಟ್ರ್ಯಾಕ್ಷನ್‌ಗಳ ಕೊರತೆಯಿಂದಾಗಿ ಕರ್ನೆಲ್‌ಗಳು ಸ್ವಲ್ಪ ಹೆಚ್ಚು ವರ್ಬೋಸ್ ಆಗಿದ್ದರೂ, ಹಾರ್ಡ್‌ವೇರ್ ಮೇಲೆ ಉತ್ತಮ-ಗ್ರಿನ್ಯುಲಾರಿಟಿ ನಿಯಂತ್ರಣವು ವೇಗವಾದ ಕರ್ನೆಲ್‌ಗಳನ್ನು ನಿರ್ಮಿಸಲು ನಮಗೆ ಅನುಮತಿಸುವುದರಿಂದ ಸಂಕೀರ್ಣತೆಯು ಸಮರ್ಥಿಸಲ್ಪಟ್ಟಿದೆ ಎಂದು ನಾವು ಕಂಡುಕೊಂಡಿದ್ದೇವೆ.

ಬರೆಯುವ ಸಮಯದಲ್ಲಿ, CUTLASS C++ ಮತ್ತು CuTeDSL ಕರ್ನೆಲ್‌ಗಳ ನಡುವಿನ ಕಾರ್ಯಕ್ಷಮತೆ ಅಂತರವು ಕಡಿಮೆಯಾಗಿದೆ, ಇದು ಸಾಧನದ ನಡುವಿನ ಸಂವಹನ, ನಾರ್ಮಲೈಸೇಶನ್ ಅಥವಾ ಆಕ್ಟಿವೇಷನ್‌ಗಳೊಂದಿಗೆ ಹೆಚ್ಚಿನ ಕಾರ್ಯಕ್ಷಮತೆಯ GEMM ಕರ್ನೆಲ್‌ಗಳನ್ನು ಸಂಯೋಜಿಸಲು ಮತ್ತು ಫ್ಯೂಸ್ ಮಾಡಲು ರೋಮಾಂಚಕಾರಿ ಅವಕಾಶಗಳನ್ನು ತೆರೆಯುತ್ತದೆ. ಈ DSL ಗೆ ಬದ್ಧರಾಗಿರುವ ಮೂಲಕ, ಎಂಭೆಡ್ಡಿಂಗ್‌ಗಳಿಂದ ಹಿಡಿದು LLM ಗಳವರೆಗಿನ ನಮ್ಮ ಕೆಲಸದ ಹೊರೆಗಳಾದ್ಯಂತ ಥ್ರುಪುಟ್ ಅನ್ನು ಸುಧಾರಿಸಲು ನಮ್ಮ ಸಾಧನಗಳಿಂದ ಹೆಚ್ಚಿನ ಕಾರ್ಯಕ್ಷಮತೆಯನ್ನು ಹೊರತೆಗೆಯಲು ನಮಗೆ ಸಾಧ್ಯವಾಗುತ್ತದೆ.

https://github.com/Dao-AILab/quack

https://github.com/NVIDIA/cutlass

https://github.com/NVIDIA/nccl-tests

https://docs.nvidia.com/cuda/parallel-thread-execution/

BibTeX

@online{perplexity2026CutedslAtPerplexity,
  author = {{Perplexity} and Nandor Licker and Bex Brooker and Jesse Dwyer and Phi Hoang},
  title = {{CuTeDSL at Perplexity}},
  date = {2026-05-06},
  url = {https://www.perplexity.ai/hub/blog/cutedsl-at-perplexity}
}

Perplexity AI ನಲ್ಲಿ ಇನ್‌ಫೆರೆನ್ಸ್

CuTeDSL ಏಕೆ?

ಡೆವಲಪರ್ ಅನುಭವ

ದೊಡ್ಡ ಕಾನ್ಫಿಗರೇಶನ್ ಸ್ಪೇಸ್

ಪ್ರೀಫಿಲ್ ಮತ್ತು ಡಿಕೋಡ್ ವಿಶೇಷತೆ

ಗ್ರಿಡ್ ಸಿಂಕ್ರೊನೈಸೇಶನ್

CuTeDSL ಕರ್ನೆಲ್‌ಗಳು

QK ನಾರ್ಮ್

ಮಿಕ್ಸ್ಚರ್-ಆಫ್-ಎಕ್ಸ್‌ಪರ್ಟ್ಸ್ ಡಿಸ್ಪ್ಯಾಚ್/ಕಂಬೈನ್

ಸ್ಥಳೀಯ

ಸಿಂಗಲ್-ನೋಡ್

ಡೇಟಾ ಪ್ಯಾರಲಲ್

ಪೀರ್-ಟು-ಪೀರ್

Blackwell ಬೆಂಬಲ

AI-ನೆರವಿನ ಕರ್ನೆಲ್ ಉತ್ಪಾದನೆ

ತೀರ್ಮಾನ ಮತ್ತು ಭವಿಷ್ಯದ ಕೆಲಸ

ಉಲ್ಲೇಖಗಳು

ಉಲ್ಲೇಖ