The most general kernel launch we can do looks like thi:, square of 3 parameters.
The first is the dimensionality of the grid of blocks
that has bx X by X bz blocks.
Each one of those blocks is specified by this parameter: the block of threads that has tx X ty X tz threads in it,
and recall that this has a maximum size.
Finally, there's a third argument that defaults to zero if you don't use it,
and we're not going to cover it specifically today.
It's the amount of shared memory in bytes allocated per thread block.
With this one kernel call, you can launch an enormous number of threads.
And let's all remember, with great power comes great responsibility, so launch your kernels wisely.
One more important thing about blocks and threads.
Recall from our square kernel, that each thread knows its thread ID within a block.
It actually knows many things.
First is threaded x, as we've seen, which thread it is within the block.
Here we have a block.
Each thread, say this thread here, knows its index in each of the x, y, and z dimensions,
and we can access those as thread idx.x, thread idx.y, and dot z.
We also know block Dim, the size of a block.
How many threads are there in this block
along the x dimension, the y dimension, and potentially the z dimension?
So we know those two things for a block.
We know the analogous things for a grid.
Block index for instance is which block am I in within the grid. Again dot x, dot y, and dot z.
And grid Dim will tell us the size of the grid, how many blocks there are
in the x dimension, the y dimension, and the z dimension.
What I want you to take home from this little discussion is only the following.
It's convenient to have multi-dimensional grids and blocks when your problem has multiple dimensions.
CUDA implements this natively and efficiently.
When you call thread at idx.x, or block dim.y, that's a very efficient thing within CUDA.
Since we're doing image processing in this course,
you should be counting on finding a lot of two dimensional grids and blocks.
So, let's wrap up with a little quiz.
Let's say I launch the following kernel.
Kernel with 2 parameters dim 3 (8, 4, 2, 2) and dim 3 (16, 16).
How many blocks will this call launch,
how many threads per block, and how many total threads?
Наиболее общая форма запуска kernel выглядит так, square от трех параметров.
Первый размерность сетки блоков
здесь bx x by x bz блоков.
Каждый из блоков задан таким параметром: блок thread'ов по tx x ty x tz в каждом,
и помним, что у них есть максимальный размер.
Наконец, есть третий параметр, который по умолчанию 0, если мы его не используем,
и сегодня мы не будем его рассматривать.
Это количества shared памяти в байтах аллоцированной для каждого блока.
С таким вызовом kernel'а вы можете запускать огромное количество thread'ов.
Только помним, что в великой силой приходит великая ответственность, т.е. запускаем kernel'ы разумно.
Еще одна важная вещь о блоках и thread'ах.
Вспомним из нашего kernel'а square, что каждый thread знает свой thread ID в блоке.
Он на самом деле многое знает.
Во-первых это threadID.x, как мы уже видели это определяет что за поток в блоке.
Здесь у нас блок.
Каждый thread, например этот, знает свой индекс в каждом из измерений x, y, z,
и мы можем получить к ним доступ как threadId.x, threadId.y, threadId.z.
Мы также знаем blockDim, размер блока.
Т.е. Как много thread'ов в каждом блоке
по направлению x, y и z.
Итак мы знаем это про блок.
Аналогичной мы знаем для сетки.
blockIdx, например, определяет что это за блок в сетке. И опять c .x, .y, .z.
И gridDim говорит нам о размере сетки, как много блоков в ней
по направлению x, направлению y, и направлению z.
Я бы хотел чтобы вы запомнили из этого разговора следующее.
Это удобно иметь многомерные сетки и блоки, когда задача имеет несколько измерений.
CUDA реализует это изначально и эффективно.
Когда вы вызываете threadId.x или blockDim.y это очень эффективно в CUDA.
Т.к. у нас будет много обработки изображений в это курсе,
вам следует рассчитывать на обнаружение множества двумерных сеток и блоков.
Завершим небольшим опросом.
Пусть я запускаю следующее ядро.
С двумя параметрами dim3(8, 4, 2, 2) и размерностью dim3(16, 16).
Как много блоков запустится при этом вызове,
как много thread'ов в блоке и какое их итоговое кол-во.
我们能做的最普通的内核启动看起来像这个:
3 个参数的平方。
第一个参数是网格块的维数,
那有 bx X by X bz 个块。
这些块中的每一个都由这个参数指定:
具有 tx X ty X tz 个线程的线程块,
回顾一下,这有一个最大尺寸。
最后,还有第三个参数,默认值为零,如果你不使用它,
我们今天不会专门讨论它。
它是以字节表示的每个线程块分配的共享内存量。
借助这一个内核调用,你可以启动大量线程。
让我们都记住,能力越大,责任越大,所以明智地启动你的内核。
另一件有关块和线程的重要事情。
从我们的平方内核回忆一下,
每个线程知道其在一个块内的线程 ID。
它其实知道很多事情。
如我们所见,第一个是线程id x,它是在块内的线程。
在这里我们有一个块。
每个线程,比如这里的这个线程,
知道其索引中的每个 x、y和z的维度值,
那么我们能以线程idx.x、线程idx.y和线程idx.z来访问那些线程。
我们也知道了块的维度,即块的大小。
这个块中有多少线程
在x 方向,y 方向,还有可能的 z 方向?
所以我们知道了块的这两件事情。
我们知道网格的类似事情。
例如,块索引,就是我在网格内的哪一块。同理,点x、 点y和点z。
网格维度会告诉我们网格的大小,有多少块位于
x 方向、 y 方向和 z 方向。
我需要你从这小讨论中明白的只是以下内容。
当你的问题有多个维度时,拥有多维网格和块很方便。
CUDA与生俱来能够高效地实现这个。
当您调用 idx.x,或块 dim.y,这在CUDA中非常高效。
因为我们在这课程中做图像处理,
你应该指望找到很多的二维网格和块。
让我们以小测验来结束。
比方说我启动了以下内核。
带有 dim 3(8, 4, 2,2)和dim 3 (16,16)两个参数的内核。
这个调用将启动多少块,
每个块有多少线程,另外总的线程有多少?