Какой самый эффективный способ вычислить идентификатор деформации / линии в 1-D сетке?

В CUDA каждый поток знает свой индекс блока в сетке и индекс потока в блоке. Но два важных значения явно не доступны для него:

  • Индекс в качестве полосы в пределах его варпа (его "идентификатор полосы")
  • Индекс деформации, для которой он является полосой движения внутри блока (его "идентификатор деформации")

Предполагая, что сетка является одномерной (она же линейная, т.е. blockDim.y а также blockDim.z 1), очевидно, можно получить их следующим образом:

enum : unsigned { warp_size = 32 };
auto lane_id = threadIdx.x % warp_size;
auto warp_id = threadIdx.x / warp_size;

и если вы не доверяете компилятору оптимизировать это, вы можете переписать его так:

enum : unsigned { warp_size = 32, log_warp_size = 5 };
auto lane_id = threadIdx.x & (warp_size - 1);
auto warp_id = threadIdx.x >> log_warp_size;

это самая эффективная вещь, чтобы сделать? Кажется, что все потоки приходится тратить на это много времени.

(вдохновлен этим вопросом.)

Источник

2 ответа

Наивные вычисления в настоящее время наиболее эффективны.

Примечание: этот ответ был сильно отредактирован.

Очень заманчиво попытаться вообще избежать вычислений - так как эти два значения, кажется, уже доступны, если вы загляните под капот.

Видите ли, графические процессоры nVIDIA имеют специальные регистры, которые ваш (скомпилированный) код может читать для доступа к различным видам полезной информации. Один такой регистр содержит threadIdx.x; другой держит blockDim.x; другой - счетчик тактов; и так далее. C++ как язык не раскрывает их, очевидно; и, фактически, не делает CUDA. Однако промежуточное представление, в которое компилируется код CUDA, с именем PTX, действительно предоставляет эти специальные регистры (начиная с PTX 1.3, то есть с версиями CUDA>= 2.1).

Два из этих специальных регистров %warpid а также %laneid, Теперь CUDA поддерживает встраивание кода PTX в код CUDA с помощью asm ключевое слово - точно так же, как оно может быть использовано для кода на стороне хоста, чтобы напрямую генерировать инструкции сборки процессора. С этим механизмом можно использовать эти специальные регистры:

__forceinline__ __device__ unsigned lane_id()
{
    unsigned ret; 
    asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret));
    return ret;
}

__forceinline__ __device__ unsigned warp_id()
{
    // this is not equal to threadIdx.x / 32
    unsigned ret; 
    asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
    return ret;
}

... но здесь есть две проблемы.

Первая проблема, как предполагает @Patwie, заключается в том, что %warp_id не дает вам того, что вы на самом деле хотите, - это не индекс деформации в контексте сетки, а скорее в контексте физического SM (который может содержать так много варпов, проживающих одновременно), и эти два не являются тот же самый. Так что не используйте %warp_id,

Что касается %lane_id, он дает вам правильное значение, но вводит в заблуждение неэффективно: несмотря на то, что это "регистр", он не похож на обычные регистры в вашем регистре, с задержкой доступа к 1 циклу. Это специальный регистр, который в реальном оборудовании извлекается с помощью S2R инструкция, которая может проявлять большую задержку.


Итог: просто вычислите ID деформации и ID нити. Не могу обойти это - пока.

Источник

Другой ответ очень опасен! Вычислите lane-id и warp-id самостоятельно.

#include <cuda.h>
#include <iostream>

inline __device__ unsigned get_lane_id() {
  unsigned ret;
  asm volatile("mov.u32 %0, %laneid;" : "=r"(ret));
  return ret;
}

inline __device__ unsigned get_warp_id() {
  unsigned ret;
  asm volatile("mov.u32 %0, %warpid;" : "=r"(ret));
  return ret;
}

__global__ void kernel() {
  const int actual_warpid = get_warp_id();
  const int actual_laneid = get_lane_id();
  const int expected_warpid = threadIdx.x / 32;
  const int expected_laneid = threadIdx.x % 32;
  if (expected_laneid == 0) {
    printf("[warp:] actual: %i  expected: %i\n", actual_warpid,
           expected_warpid);
    printf("[lane:] actual: %i  expected: %i\n", actual_laneid,
           expected_laneid);
  }
}

int main(int argc, char const *argv[]) {
  dim3 grid(8, 7, 1);
  dim3 block(4 * 32, 1);

  kernel<<<grid, block>>>();
  cudaDeviceSynchronize();
  return 0;
}

который дает что-то вроде

[warp:] actual: 4  expected: 3
[warp:] actual: 10  expected: 0
[warp:] actual: 1  expected: 1
[warp:] actual: 12  expected: 1
[warp:] actual: 4  expected: 3
[warp:] actual: 0  expected: 0
[warp:] actual: 13  expected: 2
[warp:] actual: 12  expected: 1
[warp:] actual: 6  expected: 1
[warp:] actual: 6  expected: 1
[warp:] actual: 13  expected: 2
[warp:] actual: 10  expected: 0
[warp:] actual: 1  expected: 1
...
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0

см. также документы PTX

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

Обратите внимание, что% warpid является энергозависимым и возвращает местоположение потока в момент чтения, но его значение может измениться во время выполнения, например, из-за перепланирования потоков после выгрузки.

Следовательно, это warp-id планировщика без какой-либо гарантии того, что он совпадает с виртуальным warp-id (начинается с отсчета с 0).

Документы проясняют это:

По этой причине, %ctaid и%tid должны использоваться для вычисления индекса виртуальной деформации, если такое значение требуется в коде ядра; % warpid предназначен главным образом для того, чтобы профилирующий и диагностический код собирал и регистрировал информацию, такую ​​как отображение рабочего места и распределение нагрузки.

Если вы думаете, хорошо, давайте использовать CUB для этого: это даже влияет cub::WarpId()

Возвращает идентификатор деформации вызывающего потока. ID деформации гарантированно будет уникальным среди деформаций, но может не соответствовать ранжированию с нуля внутри блока потока.

РЕДАКТИРОВАТЬ: Использование %laneid кажется безопасным

Источник
Другие вопросы по тегам