Samlade eller stegvisa läsningar
Så påverkar mappningen mellan trådar och adresser resultatet.
Samlade eller stegvisa läsningar är en gratis lektion i CUDA Academy på CoddyKit. Detta är lektion 2 av 4. Ni kan läsa hela lektionen gratis nedan och sedan öva praktiskt i webbläsaren med en inbyggd kodredigerare och en AI-handledare som är tillgänglig dygnet runt. Den ingår i lärvägen för CUDA Academy, och Era framsteg synkroniseras mellan webben och CoddyKit-appen. Kursen i CUDA Academy innehåller totalt 4 lektioner.
Allt handlar om mappning
Sammanfogning beror på vilken adress varje tråd berör. Avbildningen från tråd till adress avgör om en transaktion betjänar hela warpen.
Det sammanfogade mönstret
När närliggande trådar läser närliggande adresser täcker warpen ett sammanhängande område. Denna prydliga layout är en sammanfogad åtkomst.
int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = data[i];Varför det lönar sig
Tråd 0 träffar index 0, tråd 1 träffar index 1 och så vidare. Alla 32 adresser ligger i en enda justerad cachelinje, så GPU:n behöver bara en transaktion.
Här kommer steglängden
En steglängd är ett fast avstånd mellan adresserna som efterföljande trådar berör. Så fort avståndet blir större än ett element börjar sammanfogningen brytas upp.
float v = data[i * stride];Stridiga läsningar sprids ut
Med en steglängd på 2 läser tråd 0 värdet 0, tråd 1 värdet 2 och tråd 2 värdet 4. Warpen täcker nu dubbelt så mycket minne och behöver därför fler transaktioner.
Kostnaden ökar
Om steglängden fördubblas fördubblas ungefär också antalet berörda cachelinjer. Stora steglängder kan göra att en warp närmar sig många transaktioner samtidigt som den bara använder en liten del av varje cachelinje.
En vanlig fälla
Om varje tråd får en hel kolumn i en radprioriterad matris skapas en enorm steglängd. Koden ser prydlig ut, men sprider i det tysta ut varje warp.
float v = matrix[threadIdx.x * width + row];Transponera avbildningen
Ofta är lösningen helt enkelt att byta vilket index som varierar snabbast med tråden. Låt efterföljande trådar gå igenom efterföljande minnesadresser, så sammanfogas läsningarna igen.
float v = matrix[row * width + threadIdx.x];Även förskjutningar ställer till det
Även ett sammanfogat mönster påverkas om det börjar mitt i en cachelinje. En feljusterad förskjutning flyttar warpen över en gräns och delar upp en läsning i två.
Tänk i warps
För att bedöma ett mönster ska Ni inte föreställa Er en enda tråd. Föreställ Er alla 32 lane-positioner samtidigt och fråga hur många cachelinjer deras adresser täcker tillsammans. 🔍
Tumregeln
Låt det snabbast föränderliga indexet följa threadIdx.x. Den enda vanan håller de flesta globala läsningarna sammanfogade utan extra arbete.
Snabb kontroll
Välj det åtkomstmönster som sammanfogas bäst.
Sammanfattning
Ni såg att sammanfogade läsningar håller närliggande värden tillsammans, medan en steglängd sprider ut dem över flera cachelinjer. Låt threadIdx.x styra det snabbast föränderliga indexet. 🎉
Lär dig C++ med en AI-lärare – gratis
Skriv och kör riktig kod i webbläsaren, få omedelbar hjälp av en AI-lärare dygnet runt och fortsätt där du slutade – på webben eller i appen.
- Kurser
- 30
- Lektioner
- 120
Vanliga frågor
Är lektionen ”Samlade eller stegvisa läsningar” gratis?
Ja – hela texten till ”Samlade eller stegvisa läsningar” kan läsas gratis här på webben. Om Ni vill öva interaktivt med en inbyggd kodredigerare och en AI-handledare som är tillgänglig dygnet runt och låsa upp resten av kursen i CUDA Academy, kan Ni uppgradera till CoddyKit PRO. Kursen i CUDA Academy innehåller totalt 4 lektioner.
Vad lär jag mig i ”Samlade eller stegvisa läsningar”?
Så påverkar mappningen mellan trådar och adresser resultatet. Ni övar på CUDA Academy med praktisk kod som körs direkt i webbläsaren, medan en AI-handledare som är tillgänglig dygnet runt svarar på Era frågor under lektionen.
Behöver jag någon erfarenhet för att börja lära mig CUDA Academy?
Du behöver inga förkunskaper. Utbildningen i CUDA Academy på CoddyKit är upplagd för allt från nybörjare till avancerade elever, så att du kan börja här eller från början och gå fram i din egen takt. Detta är lektion 2 av 4.
Hur lång tid tar lektionen ”Samlade eller stegvisa läsningar”?
De flesta CoddyKit-lektioner tar cirka 5–10 minuter. Varje lektion är kort och interaktiv, så att du gör stadiga framsteg och kan fortsätta precis där du slutade – på webben eller i appen.
Kan jag skriva och köra kod i den här CUDA Academy-lektionen?
Ja. Varje CUDA Academy-lektion innehåller en inbyggd kodredigerare, så att du kan skriva och köra riktig kod direkt i webbläsaren och få omedelbar AI-feedback – utan lokal installation.
Alla lektioner i den här kursen
- Vad en minnestransaktion är
- Samlade eller stegvisa läsningar
- Array av strukturer eller struktur av arrayer
- Mät effektiv bandbredd