Наблюдение за тем, как блоки потоков распределены для SM во время выполнения?

Я начинаю с профилированием CUDA. Я в основном хочу создать временную шкалу, которая показывает каждый SM и блок потока, который был назначен ему во время выполнения.

Что-то похожее на это:

Автор: Sreepati Pai

Я читал о чтении% smid register, но я не знаю, как включить его с кодом, который я хочу проверить, или как связать это с потоковыми блоками или временем.

2 ответа

Полный код выходит за рамки этого ответа, поэтому в этом ответе вы найдете строительные блоки для реализации блочной трассировки.

  1. Выделить буфер 16 байт * количество блоков. Это можно сделать за один запуск, или большой буфер можно выделить и сохранить для нескольких запусков.
  2. Передайте указатель блока либо через постоянную переменную, либо как дополнительный параметр ядра.
  3. Измените свои глобальные функции, чтобы принять параметр и выполнить приведенный ниже код. Я рекомендую писать новые глобальные обертки функций и иметь оболочку оболочки, вызывающую старый код. Это упрощает обработку ядер с несколькими точками выхода.

Визуализация данных

  1. В вычислительной способности 2.x устройства функция timestamp должна быть clock64. Эти часы не синхронизируются через SM. Рекомендуемый подход состоит в том, чтобы сортировать время на SM и использовать самое низкое время для SM как время запуска ядра. Это будет отключено только на 100 секунд от реального времени, поэтому для ядер разумного размера этот дрейф пренебрежимо мал.
  2. Удалите smid из нижних 4-бит первого 8-байтового значения. Очистите нижние 4-бит временной метки времени.

Выделите буфер устройства, равный количеству блоков * 16 байтов. Каждые 16-байтовые записи будут хранить начальную и конечную временную метку, а также 5-битное smid-приложение, упакованное во время начала.

static __device__ inline ******** __smid()
{
 ******** smid;
 asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
 return smid;
}

// use globaltimer for compute capability >= 3.0 (kepler and maxwell)
// use clock64 for compute capability 2.x (fermi)
static __device__ inline ******** __timestamp()
{
 ******** globaltime;
 asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(globaltime) );
 return globaltime;
}

__global__ blocktime(********* pBlockTime)
{
 // START TIMESTAMP
 ******** startTime = __timestamp();
 // flatBlockIdx should be adjusted to 1D, 2D, and 3D launches to minimize
 // overhead. Reduce to ******** if launch index does not exceed 32-bit.
 ******** flatBlockIdx = (blockIdx.z * gridDim.x * gridDim.y)
 + (blockIdx.y * gridDim.x)
 + blockIdx.x;

 // reduce this based upon dimensions of block to minimize overhead
 if (threadIdx.x == 0 && theradIdx.y == 0 && threadIdx.z == 0)
 {
 // Put the smid in the 4 lower bits. If the MultiprocessCounter exceeds
 // 16 then increase to 5-bits. The lower 5-bits of globaltimer are
 // junk. If using clock64 and you want the improve precision then use
 // the most significant 4-5 bits.
 ******** smid = __smid();
 ******** data = (startTime & 0xF) | smid;
 pBlockTime[flatBlockIdx * 2 + 0] = data;
 }

 // do work

 // I would recommend changing your current __global__ function to be
 // a __global__ __device__ function and call it here. This will result
 // in easier handling of kernels that have multiple exit points.

 // END TIMESTAMP
 // All threads in block will write out. This is not very efficient.
 // Depending on the kernel this can be reduced to 1 thread or 1 thread per warp.
 ******** endTime = __timestamp();
 pBlockTime[flatBlockIdx * 2 + 1] = endTime;
}


__noinline__ __device__ **** get_smid(void)
{
 **** ret;
 asm("mov.u32 %0, %smid;" : "=r"(ret) );
 return ret;
}

Источник здесь.

licensed under cc by-sa 3.0 with attribution.