Sürücü, görevlerin ve bağımlılıkların tam tanımını kullanarak yürütmeyi optimize edebildiğinden, CUDA Grafikleri önemli bir performans artışı sağlayabilir. Grafikler, grafik oluşturma yükünün birbirini takip eden birçok başlatmayla amorti edilebildiği statik iş akışları için inanılmaz faydalar sağlar.
Bununla birlikte, neredeyse tüm problemler, grafiklerin bölünmesini ve bir sonraki hangi çalışmanın başlatılacağına karar vermek için kontrolün CPU'ya verilmesini gerektirebilen bir tür karar verme sürecini içerir. İşin bu şekilde bölünmesi, CUDA'nın optimizasyon yapma becerisini tehlikeye atar, CPU kaynaklarını bağlar ve her grafik açılışında ek yük yaratır.
CUDA 12.4'ten başlayarak CUDA Graphs şunları destekler: koşullu düğümler, kontrolü CPU'ya döndürmeden bir grafiğin bölümlerinin koşullu veya tekrarlı yürütülmesini sağlar. Bu, CPU kaynaklarını serbest bırakarak çok daha fazla iş akışının tek bir grafikte temsil edilmesini sağlar.
Koşullu düğümler
Koşullu düğümlerin iki çeşidi vardır:
- IF düğümleri: Koşul değeri doğruysa, düğüm her değerlendirildiğinde gövde bir kez yürütülür.
- WHILE düğümleri: Koşul değeri doğru olduğu sürece düğüm değerlendirildiğinde gövde tekrar tekrar yürütülür.
Koşullu düğümler, alt grafik düğümlerine benzer şekilde konteyner düğümleridir, ancak düğüm içinde yer alan grafiğin yürütülmesi, bir koşul değişkeninin değerine bağlıdır. Bir düğümle ilişkili koşul değerine, düğümden önce oluşturulması gereken bir tanıtıcı tarafından erişilir. Koşul değeri bir CUDA çekirdeğinde çağrılarak ayarlanabilir. cudaGraphSetConditional
. Grafiğin her başlangıcında uygulanan bir başlatma, tanıtıcı oluşturulduğunda da belirtilebilir.
Koşullu düğüm oluşturulduğunda boş bir grafik de oluşturulur ve tanıtıcı kullanıcıya döndürülür. Bu grafik düğüme bağlanır ve koşul değerine göre yürütülür. Bu koşullu gövde grafiği aşağıdaki yöntemlerden biri kullanılarak doldurulabilir: grafik API'si veya kullanarak eşzamansız CUDA çağrılarını yakalayarak cudaStreamBeginCaptureToGraph
.
Koşullu düğümler aynı zamanda iç içe de yerleştirilebilir. Örneğin, koşullu IF düğümü içeren bir gövde grafiğiyle koşullu bir WHILE düğümü oluşturabilirsiniz.
Koşullu düğüm gövdesi grafikleri aşağıdakilerden herhangi birini içerebilir:
- Çekirdek düğümleri (CNP, kooperatif şu anda desteklenmiyor)
- Boş düğümler
- Alt grafik düğümleri
- Memset düğümleri
- Bellek kopyası düğümleri
- Koşullu düğümler
Bu, alt grafiklere ve koşullu gövdelere yinelemeli olarak uygulanır. İç içe koşullu ifadelerdeki veya herhangi bir düzeydeki alt grafiklerdeki çekirdekler de dahil olmak üzere tüm çekirdekler, aynı CUDA bağlamına ait olmalıdır. Bellek kopyaları ve bellek kümeleri, koşullu düğümün bağlamından erişilebilen bellek üzerinde hareket etmelidir.
Tam örnekler mevcuttur CUDA numune deposu. Bir sonraki bölümde, koşullu düğümlerle neler yapabileceğinizi göstermek için bazı örnekler üzerinden geçilecektir.
Koşullu IF düğümleri
Bir IF düğümünün gövde grafiği, IF düğümü her değerlendirildiğinde koşulun sıfırdan farklı olması durumunda bir kez yürütülür. Şekil 1, orta düğüm B'nin dört düğümlü bir grafik içeren bir IF koşullu düğüm olduğu bir grafiği göstermektedir:
Bu grafiğin nasıl oluşturulabileceğini göstermek için aşağıdaki örnek, koşullu düğümün (B) yukarı akışındaki bir çekirdek olan A düğümünü kullanarak, o çekirdek tarafından yapılan işin sonuçlarına dayalı olarak koşullunun değerini ayarlar. Koşulun gövdesi grafik API'si kullanılarak doldurulur.
İlk olarak, A düğümünün çekirdeğini tanımlayın. Bu çekirdek, kullanıcı tarafından gerçekleştirilen bazı keyfi hesaplamaların sonucuna bağlı olarak koşullu tanıtıcıyı ayarlar.
__global__ void setHandle(cudaGraphConditionalHandle handle)
{
unsigned int value = 0;
// We could perform some work here and set value based on the result of that work.
if (someCondition) {
// Set ‘value’ to non-zero if we want the conditional body to execute
value = 1;
}
cudaGraphSetConditional(handle, value);
}
Daha sonra grafiği oluşturmak için bir fonksiyon tanımlayın. Bu işlev koşullu tanıtıcıyı tahsis eder, düğümleri oluşturur ve koşullu grafiğin gövdesini doldurur. Anlaşılır olması açısından grafiği başlatacak ve yürütecek kod atlanmıştır.
cudaGraph_t createGraph() {
cudaGraph_t graph;
cudaGraphNode_t node;
void *kernelArgs[1];
cudaGraphCreate(&graph, 0);
cudaGraphConditionalHandle handle;
cudaGraphConditionalHandleCreate(&handle, graph);
// Use a kernel upstream of the conditional to set the handle value
cudaGraphNodeParams kParams = { cudaGraphNodeTypeKernel };
kParams.kernel.func = (void *)setHandle;
kParams.kernel.gridDim.x = kParams.kernel.gridDim.y = kParams.kernel.gridDim.z = 1;
kParams.kernel.blockDim.x = kParams.kernel.blockDim.y = kParams.kernel.blockDim.z = 1;
kParams.kernel.kernelParams = kernelArgs;
kernelArgs[0] = &handle;
cudaGraphAddNode(&node, graph, NULL, 0, &kParams);
cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional };
cParams.conditional.handle = handle;
cParams.conditional.type = cudaGraphCondTypeIf;
cParams.conditional.size = 1;
cudaGraphAddNode(&node, graph, &node, 1, &cParams);
cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0];
// Populate the body of the conditional node
cudaGraphNode_t bodyNodes[4];
cudaGraphNodeParams params[4] = { ... }; // Setup kernel parameters as needed.
cudaGraphAddNode(&bodyNodes[0], bodyGraph, NULL, 0, ¶ms[0]);
cudaGraphAddNode(&bodyNodes[1], bodyGraph, &bodyNodes[0], 1, ¶ms[1]);
cudaGraphAddNode(&bodyNodes[2], bodyGraph, &bodyNodes[0], 1, ¶ms[2]);
cudaGraphAddNode(&bodyNodes[3], bodyGraph, &bodyNodes[1], 2, ¶ms[3]);
return graph;
}
Koşullu WHILE düğümleri
WHILE düğümünün gövde grafiği, koşul sıfırdan farklı olduğu sürece tekrar tekrar yürütülür. Koşul, düğüm yürütüldüğünde ve gövde grafiğinin her tamamlanmasından sonra değerlendirilecektir. Aşağıdaki diyagramda ortadaki düğüm B'nin üç düğümlü bir grafik içeren WHILE koşullu düğüm olduğu üç düğümlü bir grafik gösterilmektedir.
Bu grafiğin nasıl oluşturulabileceğini görmek için aşağıdaki örnek, tanıtıcının varsayılan değerini sıfırdan farklı bir değere ayarlar, böylece WHILE döngüsü varsayılan olarak yürütülür. Varsayılan değeri sıfırdan farklı bir değere ayarlamak ve koşullu değeri, koşullu değiştirilmemiş çekirdeğin yukarı akışında bırakmak, koşullu gövdenin her zaman en az bir kez çalıştırıldığı bir do-while döngüsü oluşturur. Döngü gövdesinin yalnızca koşul doğru olduğunda yürütüldüğü bir WHILE döngüsü oluşturmak, bazı hesaplamaların yapılmasını ve A düğümünde koşul tutamacının uygun şekilde ayarlanmasını gerektirir.
Önceki örnekte koşullu gövde grafik API'si ile doldurulmuştur. Bu örnekte koşulun gövdesi akış yakalama kullanılarak doldurulur.
İlk adım, koşullu gövdenin her yürütülmesi sırasında koşullu değeri ayarlayan bir çekirdek tanımlamaktır. Bu örnekte tutamaç, aşağı sayacın değerine göre ayarlanmıştır.
__global__ void loopKernel(cudaGraphConditionalHandle handle)
{
static int count = 10;
cudaGraphSetConditional(handle, --count ? 1 : 0);
}
Daha sonra grafiği oluşturmak için bir fonksiyon tanımlayın. Bu işlev koşullu tanıtıcıyı tahsis eder, düğümleri oluşturur ve koşullu grafiğin gövdesini doldurur. Anlaşılır olması açısından grafiği başlatacak ve yürütecek kod atlanmıştır.
cudaGraph_t createGraph() {
cudaGraph_t graph;
cudaGraphNode_t nodes[3];
cudaGraphCreate(&graph, 0);
// Insert kernel node A
cudaGraphNodeParams params = ...;
cudaGraphAddNode(&nodes[0], graph, NULL, 0, ¶ms);
cudaGraphConditionalHandle handle;
cudaGraphConditionalHandleCreate(&handle, graph, 1, cudaGraphCondAssignDefault);
// Insert conditional node B
cudaGraphNodeParams cParams = { cudaGraphNodeTypeConditional };
cParams.conditional.handle = handle;
cParams.conditional.type = cudaGraphCondTypeWhile;
cParams.conditional.size = 1;
cudaGraphAddNode(&nodes[1], graph, &nodes[0], 1, &cParams);
cudaGraph_t bodyGraph = cParams.conditional.phGraph_out[0];
cudaStream_t captureStream;
cudaStreamCreate(&captureStream);
// Fill out body graph with stream capture.
cudaStreamBeginCaptureToGraph(captureStream,
bodyGraph,
nullptr,
nullptr,
0,
cudaStreamCaptureModeRelaxed);
myKernel1<<<..., captureStream>>>(...);
myKernel2<<<..., captureStream>>>(...);
loopKernel<<<1, 1, 0, captureStream>>>(handle);
cudaStreamEndCapture(captureStream, nullptr);
cudaStreamDestroy(captureStream);
// Insert kernel node C.
params = ...;
cudaGraphAddNode(&nodes[2], graph, &nodes[1], 1, ¶ms);
return graph;
}
Bu örnek şunu kullanır: cudaStreamBeginCaptureToGraph
CUDA 12.3'e eklenen yeni bir API olup, mevcut bir grafiğe düğüm eklemek için akış yakalamayı sağlar. Bu API kullanılarak birden çok ayrı yakalama tek bir grafik nesnesinde birleştirilebilir. Bu API aynı zamanda koşullu düğümle birlikte oluşturulan koşullu gövde grafiği nesnesinin doldurulmasına da olanak tanır.
Çözüm
CUDA Grafikleri, grafik oluşturma yükünün birbirini takip eden birçok başlatmayla amorti edilebildiği statik iş akışları için inanılmaz faydalar sağlar. Grafikleri bölmeyi ortadan kaldırmak ve hangi başlatmanın önceliklendirileceğine karar vermek için kontrolü CPU'ya vermek, CPU yükünü ve gecikmeyi azaltmaya yardımcı olur. CUDA Grafiklerini koşullu düğümlerle kullanmak, kontrolü CPU'ya döndürmeden bir grafiğin bölümlerinin koşullu veya tekrarlı olarak yürütülmesini sağlar. Bu, CPU kaynaklarını serbest bırakır ve tek bir grafiğin önemli ölçüde daha karmaşık iş akışlarını temsil etmesine olanak tanır.
Koşullu düğümler hakkında daha fazla bilgi için bkz. CUDA Programlama Kılavuzu. Basit, eksiksiz örnekleri keşfetmek için şu adresi ziyaret edin: NVIDIA/cuda örnekleri GitHub'da. Ve NVIDIA Geliştirici CUDA forumlarındaki sohbete katılın.