Отчет об использовании регистра NVCC в функции __device__

Я пытаюсь получить некоторую информацию об использовании регистра в моих ядрах CUDA, используя опцию NVCC
--ptxas-options=v и хотя с глобальными функциями все в порядке, у меня возникли некоторые трудности с устройством, так как

ptxas info : Used N registers

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

У вас есть какие-нибудь предложения?

Спасибо!

Источник

1 ответ

Как я понимаю, ptxas (ассемблер устройства) выводит регистр только по коду, который он связывает. Standalone __device__ функции не связаны ассемблером, они только скомпилированы. Следовательно, ассемблер не будет выдавать значение счетчика для функций устройства. Я не верю, что есть обходной путь для этого.

Тем не менее, все еще возможно получить регистр __device__ функция путем сброса данных эльфа из вывода ассемблера, используя cuobjdump, Вы можете сделать это следующим образом:

$ cat vdot.cu
__device__  __noinline__ float vdot(float v1, float v2) {
    return (v1 * v2);
}

__device__ __noinline__  float vdot(float2 v1, float2 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y);
}

__device__ __noinline__ float vdot(float4 v1, float4 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w);
}

$ nvcc -std=c++11 -arch=sm_52 -dc -Xptxas="-v" vdot.cu
ptxas info    : 0 bytes gmem
ptxas info    : Function properties for cudaDeviceGetAttribute
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdotff
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdot6float4S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaMalloc
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaGetDevice
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdot6float2S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaFuncGetAttributes
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads

Здесь у нас есть отдельно скомпилированный набор из трех __device__ функции в объектном файле устройства. Бег cuobjdump на нем будет много выходных данных, но в нем вы получите количество регистров для каждой функции:

$ cuobjdump -elf ./vdot.o

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
compressed

<---Snipped--->


.text._Z4vdotff
bar = 0 reg = 6 lmem=0  smem=0
0xfec007f1  0x001fc000  0x00570003  0x5c980780  
0x00470000  0x5c980780  0x00370004  0x5c680000  
0xffe007ff  0x001f8000  0x0007000f  0xe3200000  
0xff87000f  0xe2400fff  0x00070f00  0x50b00000

Во второй строке вывода для функции устройства dot(float, float) Вы можете видеть, что функция использует 6 регистров. Это единственный способ изучить следы регистра функций устройства.

Источник

Я не знаю, когда он был добавлен, но моя CUDA 10 cuobjdump имеет -res-usage флаг, который показывает что-то вроде этого:

$ cuobjdump -res-usage .../cuda_compile_1_generated_VisualOdometry.cu.o

Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
identifier = /home/mad/automy-system/vision/src/VisualOdometry.cu

Resource usage:
 Common:
  GLOBAL:0 CONSTANT[3]:24
 Function _Z17vo_compute_systemPfS_P6float4S_jS0_S0_f:
  REG:39 STACK:32 SHARED:168 LOCAL:0 CONSTANT[0]:404 CONSTANT[2]:80 TEXTURE:0 SURFACE:0 SAMPLER:0
 Function _Z13vo_pre_filterP6float4PfPjPK5uint2iijff:
  REG:16 STACK:0 SHARED:8 LOCAL:0 CONSTANT[0]:372 TEXTURE:0 SURFACE:0 SAMPLER:0
Источник
Другие вопросы по тегам