Объединение памяти и результаты nvprof на NVIDIA Pascal - PullRequest
0 голосов
/ 15 мая 2019

Я провожу эксперимент по объединению памяти на Паскале и получаю неожиданные nvprof результаты. У меня есть одно ядро, которое копирует 4 ГБ поплавков из одного массива в другой. nvprof сообщает о сбивающих с толку числах для gld_transactions_per_request и gst_transactions_per_request.

Я провел эксперимент на TITAN Xp и GeForce GTX 1080 TI. Те же результаты.

#include <stdio.h>
#include <cstdint>
#include <assert.h>

#define N 1ULL1024*1024

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}


__global__ void copy_kernel(
      const float* __restrict__ data, float* __restrict__ data2) {
  for (unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
       tid < N; tid += blockDim.x * gridDim.x) {
    data2[tid] = data[tid];
  }
}

int main() {
  float* d_data;
  gpuErrchk(cudaMalloc(&d_data, sizeof(float) * N));
  assert(d_data != nullptr);
  uintptr_t d = reinterpret_cast<uintptr_t>(d_data);
  assert(d%128 == 0);  // check alignment, just to be sure

  float* d_data2;
  gpuErrchk(cudaMalloc(&d_data2, sizeof(float)*N));
  assert(d_data2 != nullptr);

  copy_kernel<<<1024,1024>>>(d_data, d_data2);
  gpuErrchk(cudaDeviceSynchronize());
}

Скомпилировано с CUDA версии 10.1:

nvcc coalescing.cu -std=c++11 -Xptxas -dlcm=ca -gencode arch=compute_61,code=sm_61 -O3

Профилируется с:

nvprof -m all ./a.out

В результатах профилирования есть несколько запутанных частей:

  • gld_transactions = 536870914, что означает, что каждая транзакция глобальной нагрузки должна быть в среднем 4GB/536870914 = 8 bytes. Это согласуется с gld_transactions_per_request = 16.000000: каждая деформация считывает 128 байтов (1 запрос), и если каждая транзакция составляет 8 байтов, то нам нужно 128 / 8 = 16 транзакций на запрос. Почему это значение так низко? Я ожидал бы идеального объединения, поэтому что-то вроде 4 (или даже 1) транзакций / запроса.
  • Запрошенная и достигнутая глобальная пропускная способность загрузки / хранения (gld_requested_throughput, gst_requested_throughput, gld_throughput, gst_throughput) составляет 150.32GB/s каждая. Я ожидаю более низкой пропускной способности для нагрузок, чем для магазинов, поскольку у нас больше транзакций на запрос.
  • gld_transactions = 536870914, но l2_read_transactions = 134218800. Глобальная память всегда доступна через кэш L1 / L2. Почему количество транзакций чтения L2 намного меньше? Это не может быть все кэшировано в L1. (global_hit_rate = 0%)

Я думаю, что неправильно прочитал nvprof результаты. Любые предложения будут оценены.

Вот полный результат профилирования:

Device "GeForce GTX 1080 Ti (0)"
    Kernel: copy_kernel(float const *, float*)
          1                             inst_per_warp                                                 Instructions per warp  1.4346e+04  1.4346e+04  1.4346e+04
          1                         branch_efficiency                                                     Branch Efficiency     100.00%     100.00%     100.00%
          1                 warp_execution_efficiency                                             Warp Execution Efficiency     100.00%     100.00%     100.00%
          1         warp_nonpred_execution_efficiency                              Warp Non-Predicated Execution Efficiency      99.99%      99.99%      99.99%
          1                      inst_replay_overhead                                           Instruction Replay Overhead    0.000178    0.000178    0.000178
          1      shared_load_transactions_per_request                           Shared Memory Load Transactions Per Request    0.000000    0.000000    0.000000
          1     shared_store_transactions_per_request                          Shared Memory Store Transactions Per Request    0.000000    0.000000    0.000000
          1       local_load_transactions_per_request                            Local Memory Load Transactions Per Request    0.000000    0.000000    0.000000
          1      local_store_transactions_per_request                           Local Memory Store Transactions Per Request    0.000000    0.000000    0.000000
          1              gld_transactions_per_request                                  Global Load Transactions Per Request   16.000000   16.000000   16.000000
          1              gst_transactions_per_request                                 Global Store Transactions Per Request    4.000000    4.000000    4.000000
          1                 shared_store_transactions                                             Shared Store Transactions           0           0           0
          1                  shared_load_transactions                                              Shared Load Transactions           0           0           0
          1                   local_load_transactions                                               Local Load Transactions           0           0           0
          1                  local_store_transactions                                              Local Store Transactions           0           0           0
          1                          gld_transactions                                              Global Load Transactions   536870914   536870914   536870914
          1                          gst_transactions                                             Global Store Transactions   134217728   134217728   134217728
          1                  sysmem_read_transactions                                       System Memory Read Transactions           0           0           0
          1                 sysmem_write_transactions                                      System Memory Write Transactions           5           5           5
          1                      l2_read_transactions                                                  L2 Read Transactions   134218800   134218800   134218800
          1                     l2_write_transactions                                                 L2 Write Transactions   134217741   134217741   134217741
          1                           global_hit_rate                                     Global Hit Rate in unified l1/tex       0.00%       0.00%       0.00%
          1                            local_hit_rate                                                        Local Hit Rate       0.00%       0.00%       0.00%
          1                  gld_requested_throughput                                      Requested Global Load Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                  gst_requested_throughput                                     Requested Global Store Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                            gld_throughput                                                Global Load Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                            gst_throughput                                               Global Store Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                     local_memory_overhead                                                 Local Memory Overhead       0.00%       0.00%       0.00%
          1                        tex_cache_hit_rate                                                Unified Cache Hit Rate      50.00%      50.00%      50.00%
          1                      l2_tex_read_hit_rate                                           L2 Hit Rate (Texture Reads)       0.00%       0.00%       0.00%
          1                     l2_tex_write_hit_rate                                          L2 Hit Rate (Texture Writes)       0.00%       0.00%       0.00%
          1                      tex_cache_throughput                                              Unified Cache Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                    l2_tex_read_throughput                                         L2 Throughput (Texture Reads)  150.32GB/s  150.32GB/s  150.32GB/s
          1                   l2_tex_write_throughput                                        L2 Throughput (Texture Writes)  150.32GB/s  150.32GB/s  150.32GB/s
          1                        l2_read_throughput                                                 L2 Throughput (Reads)  150.32GB/s  150.32GB/s  150.32GB/s
          1                       l2_write_throughput                                                L2 Throughput (Writes)  150.32GB/s  150.32GB/s  150.32GB/s
          1                    sysmem_read_throughput                                         System Memory Read Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                   sysmem_write_throughput                                        System Memory Write Throughput  5.8711KB/s  5.8711KB/s  5.8701KB/s
          1                     local_load_throughput                                          Local Memory Load Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                    local_store_throughput                                         Local Memory Store Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                    shared_load_throughput                                         Shared Memory Load Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                   shared_store_throughput                                        Shared Memory Store Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                            gld_efficiency                                         Global Memory Load Efficiency     100.00%     100.00%     100.00%
          1                            gst_efficiency                                        Global Memory Store Efficiency     100.00%     100.00%     100.00%
          1                    tex_cache_transactions                                            Unified Cache Transactions   134217728   134217728   134217728
          1                             flop_count_dp                           Floating Point Operations(Double Precision)           0           0           0
          1                         flop_count_dp_add                       Floating Point Operations(Double Precision Add)           0           0           0
          1                         flop_count_dp_fma                       Floating Point Operations(Double Precision FMA)           0           0           0
          1                         flop_count_dp_mul                       Floating Point Operations(Double Precision Mul)           0           0           0
          1                             flop_count_sp                           Floating Point Operations(Single Precision)           0           0           0
          1                         flop_count_sp_add                       Floating Point Operations(Single Precision Add)           0           0           0
          1                         flop_count_sp_fma                       Floating Point Operations(Single Precision FMA)           0           0           0
          1                         flop_count_sp_mul                        Floating Point Operation(Single Precision Mul)           0           0           0
          1                     flop_count_sp_special                   Floating Point Operations(Single Precision Special)           0           0           0
          1                             inst_executed                                                 Instructions Executed   470089728   470089728   470089728
          1                               inst_issued                                                   Instructions Issued   470173430   470173430   470173430
          1                        sysmem_utilization                                             System Memory Utilization     Low (1)     Low (1)     Low (1)
          1                          stall_inst_fetch                              Issue Stall Reasons (Instructions Fetch)       0.79%       0.79%       0.79%
          1                     stall_exec_dependency                            Issue Stall Reasons (Execution Dependency)       1.46%       1.46%       1.46%
          1                   stall_memory_dependency                                    Issue Stall Reasons (Data Request)      96.16%      96.16%      96.16%
          1                             stall_texture                                         Issue Stall Reasons (Texture)       0.00%       0.00%       0.00%
          1                                stall_sync                                 Issue Stall Reasons (Synchronization)       0.00%       0.00%       0.00%
          1                               stall_other                                           Issue Stall Reasons (Other)       1.13%       1.13%       1.13%
          1          stall_constant_memory_dependency                              Issue Stall Reasons (Immediate constant)       0.00%       0.00%       0.00%
          1                           stall_pipe_busy                                       Issue Stall Reasons (Pipe Busy)       0.07%       0.07%       0.07%
          1                         shared_efficiency                                              Shared Memory Efficiency       0.00%       0.00%       0.00%
          1                                inst_fp_32                                               FP Instructions(Single)           0           0           0
          1                                inst_fp_64                                               FP Instructions(Double)           0           0           0
          1                              inst_integer                                                  Integer Instructions  1.0742e+10  1.0742e+10  1.0742e+10
          1                          inst_bit_convert                                              Bit-Convert Instructions           0           0           0
          1                              inst_control                                             Control-Flow Instructions  1073741824  1073741824  1073741824
          1                        inst_compute_ld_st                                               Load/Store Instructions  2147483648  2147483648  2147483648
          1                                 inst_misc                                                     Misc Instructions  1077936128  1077936128  1077936128
          1           inst_inter_thread_communication                                             Inter-Thread Instructions           0           0           0
          1                               issue_slots                                                           Issue Slots   470173430   470173430   470173430
          1                                 cf_issued                                      Issued Control-Flow Instructions    33619968    33619968    33619968
          1                               cf_executed                                    Executed Control-Flow Instructions    33619968    33619968    33619968
          1                               ldst_issued                                        Issued Load/Store Instructions   268500992   268500992   268500992
          1                             ldst_executed                                      Executed Load/Store Instructions    67174400    67174400    67174400
          1                       atomic_transactions                                                   Atomic Transactions           0           0           0
          1           atomic_transactions_per_request                                       Atomic Transactions Per Request    0.000000    0.000000    0.000000
          1                      l2_atomic_throughput                                       L2 Throughput (Atomic requests)  0.00000B/s  0.00000B/s  0.00000B/s
          1                    l2_atomic_transactions                                     L2 Transactions (Atomic requests)           0           0           0
          1                  l2_tex_read_transactions                                       L2 Transactions (Texture Reads)   134217728   134217728   134217728
          1                     stall_memory_throttle                                 Issue Stall Reasons (Memory Throttle)       0.00%       0.00%       0.00%
          1                        stall_not_selected                                    Issue Stall Reasons (Not Selected)       0.39%       0.39%       0.39%
          1                 l2_tex_write_transactions                                      L2 Transactions (Texture Writes)   134217728   134217728   134217728
          1                             flop_count_hp                             Floating Point Operations(Half Precision)           0           0           0
          1                         flop_count_hp_add                         Floating Point Operations(Half Precision Add)           0           0           0
          1                         flop_count_hp_mul                          Floating Point Operation(Half Precision Mul)           0           0           0
          1                         flop_count_hp_fma                         Floating Point Operations(Half Precision FMA)           0           0           0
          1                                inst_fp_16                                                 HP Instructions(Half)           0           0           0
          1                   sysmem_read_utilization                                        System Memory Read Utilization    Idle (0)    Idle (0)    Idle (0)
          1                  sysmem_write_utilization                                       System Memory Write Utilization     Low (1)     Low (1)     Low (1)
          1               pcie_total_data_transmitted                                           PCIe Total Data Transmitted        1024        1024        1024
          1                  pcie_total_data_received                                              PCIe Total Data Received           0           0           0
          1                inst_executed_global_loads                              Warp level instructions for global loads    33554432    33554432    33554432
          1                 inst_executed_local_loads                               Warp level instructions for local loads           0           0           0
          1                inst_executed_shared_loads                              Warp level instructions for shared loads           0           0           0
          1               inst_executed_surface_loads                             Warp level instructions for surface loads           0           0           0
          1               inst_executed_global_stores                             Warp level instructions for global stores    33554432    33554432    33554432
          1                inst_executed_local_stores                              Warp level instructions for local stores           0           0           0
          1               inst_executed_shared_stores                             Warp level instructions for shared stores           0           0           0
          1              inst_executed_surface_stores                            Warp level instructions for surface stores           0           0           0
          1              inst_executed_global_atomics                  Warp level instructions for global atom and atom cas           0           0           0
          1           inst_executed_global_reductions                         Warp level instructions for global reductions           0           0           0
          1             inst_executed_surface_atomics                 Warp level instructions for surface atom and atom cas           0           0           0
          1          inst_executed_surface_reductions                        Warp level instructions for surface reductions           0           0           0
          1              inst_executed_shared_atomics                  Warp level shared instructions for atom and atom CAS           0           0           0
          1                     inst_executed_tex_ops                                   Warp level instructions for texture           0           0           0
          1                      l2_global_load_bytes       Bytes read from L2 for misses in Unified Cache for global loads  4294967296  4294967296  4294967296
          1                       l2_local_load_bytes        Bytes read from L2 for misses in Unified Cache for local loads           0           0           0
          1                     l2_surface_load_bytes      Bytes read from L2 for misses in Unified Cache for surface loads           0           0           0
          1               l2_local_global_store_bytes   Bytes written to L2 from Unified Cache for local and global stores.  4294967296  4294967296  4294967296
          1                 l2_global_reduction_bytes          Bytes written to L2 from Unified cache for global reductions           0           0           0
          1              l2_global_atomic_store_bytes             Bytes written to L2 from Unified cache for global atomics           0           0           0
          1                    l2_surface_store_bytes            Bytes written to L2 from Unified Cache for surface stores.           0           0           0
          1                l2_surface_reduction_bytes         Bytes written to L2 from Unified Cache for surface reductions           0           0           0
          1             l2_surface_atomic_store_bytes    Bytes transferred between Unified Cache and L2 for surface atomics           0           0           0
          1                      global_load_requests              Total number of global load requests from Multiprocessor   134217728   134217728   134217728
          1                       local_load_requests               Total number of local load requests from Multiprocessor           0           0           0
          1                     surface_load_requests             Total number of surface load requests from Multiprocessor           0           0           0
          1                     global_store_requests             Total number of global store requests from Multiprocessor   134217728   134217728   134217728
          1                      local_store_requests              Total number of local store requests from Multiprocessor           0           0           0
          1                    surface_store_requests            Total number of surface store requests from Multiprocessor           0           0           0
          1                    global_atomic_requests            Total number of global atomic requests from Multiprocessor           0           0           0
          1                 global_reduction_requests         Total number of global reduction requests from Multiprocessor           0           0           0
          1                   surface_atomic_requests           Total number of surface atomic requests from Multiprocessor           0           0           0
          1                surface_reduction_requests        Total number of surface reduction requests from Multiprocessor           0           0           0
          1                         sysmem_read_bytes                                              System Memory Read Bytes           0           0           0
          1                        sysmem_write_bytes                                             System Memory Write Bytes         160         160         160
          1                           l2_tex_hit_rate                                                     L2 Cache Hit Rate       0.00%       0.00%       0.00%
          1                     texture_load_requests             Total number of texture Load requests from Multiprocessor           0           0           0
          1                     unique_warps_launched                                              Number of warps launched       32768       32768       32768
          1                             sm_efficiency                                               Multiprocessor Activity      99.63%      99.63%      99.63%
          1                        achieved_occupancy                                                    Achieved Occupancy    0.986477    0.986477    0.986477
          1                                       ipc                                                          Executed IPC    0.344513    0.344513    0.344513
          1                                issued_ipc                                                            Issued IPC    0.344574    0.344574    0.344574
          1                    issue_slot_utilization                                                Issue Slot Utilization       8.61%       8.61%       8.61%
          1                  eligible_warps_per_cycle                                       Eligible Warps Per Active Cycle    0.592326    0.592326    0.592326
          1                           tex_utilization                                             Unified Cache Utilization     Low (1)     Low (1)     Low (1)
          1                            l2_utilization                                                  L2 Cache Utilization     Low (2)     Low (2)     Low (2)
          1                        shared_utilization                                             Shared Memory Utilization    Idle (0)    Idle (0)    Idle (0)
          1                       ldst_fu_utilization                                  Load/Store Function Unit Utilization     Low (1)     Low (1)     Low (1)
          1                         cf_fu_utilization                                Control-Flow Function Unit Utilization     Low (1)     Low (1)     Low (1)
          1                    special_fu_utilization                                     Special Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
          1                        tex_fu_utilization                                     Texture Function Unit Utilization     Low (1)     Low (1)     Low (1)
          1           single_precision_fu_utilization                            Single-Precision Function Unit Utilization     Low (1)     Low (1)     Low (1)
          1           double_precision_fu_utilization                            Double-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
          1                        flop_hp_efficiency                                            FLOP Efficiency(Peak Half)       0.00%       0.00%       0.00%
          1                        flop_sp_efficiency                                          FLOP Efficiency(Peak Single)       0.00%       0.00%       0.00%
          1                        flop_dp_efficiency                                          FLOP Efficiency(Peak Double)       0.00%       0.00%       0.00%
          1                    dram_read_transactions                                       Device Memory Read Transactions   134218560   134218560   134218560
          1                   dram_write_transactions                                      Device Memory Write Transactions   134176900   134176900   134176900
          1                      dram_read_throughput                                         Device Memory Read Throughput  150.32GB/s  150.32GB/s  150.32GB/s
          1                     dram_write_throughput                                        Device Memory Write Throughput  150.27GB/s  150.27GB/s  150.27GB/s
          1                          dram_utilization                                             Device Memory Utilization    High (7)    High (7)    High (7)
          1             half_precision_fu_utilization                              Half-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
          1                          ecc_transactions                                                      ECC Transactions           0           0           0
          1                            ecc_throughput                                                        ECC Throughput  0.00000B/s  0.00000B/s  0.00000B/s
          1                           dram_read_bytes                                Total bytes read from DRAM to L2 cache  4294993920  4294993920  4294993920
          1                          dram_write_bytes                             Total bytes written from L2 cache to DRAM  4293660800  4293660800  4293660800

1 Ответ

1 голос
/ 15 мая 2019

В графических процессорах Fermi и Kepler, когда выполнялась глобальная транзакция, она всегда имела размер 128 байт, а размер строки кэша L1 (если включен) составлял 128 байт.С Максвеллом и Паскалем эти характеристики изменились.В частности, чтение части кеш-линии L1 не обязательно запускает полную транзакцию шириной 128 байт.Это довольно легко обнаружить / доказать с помощью микробенчмаркинга.

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

gld_transactions = 536870914, что означает, что каждая глобальная транзакция загрузки должна быть в среднем4 ГБ / 536870914 = 8 байтов.Это согласуется с gld_transactions_per_request = 16.000000: каждая деформация считывает 128 байтов (1 запрос), и если каждая транзакция составляет 8 байтов, то нам нужно 128/8 = 16 транзакций на запрос.Почему это значение так низко?Я ожидал бы идеального слияния, поэтому что-то вроде 4 (или даже 1) транзакций / запроса.

Это склад ума (1 транзакция на запрос для полностью объединенных нагрузок 32-битного количества на поток) было бы правильно на таймфрейме Ферми / Кеплера.Это больше не подходит для графических процессоров Maxwell и Pascal.Как вы уже подсчитали, размер транзакции меньше 128 байт, и поэтому количество транзакций на запрос больше 1. Но это не указывает на проблему эффективности как таковую (как это было бы в Fermi /Кеплер сроки).Итак, давайте просто признаем, что размер транзакции может быть меньше, и поэтому транзакции на запрос могут быть выше, даже если базовый трафик по существу эффективен на 100%.

gst_transactions = 134217728 и gst_transactions_per_request = 4.000000, поэтому хранениепамять эффективнее?

Нет, это не то, что это значит.Это просто означает, что кванты подразделения могут быть разными для нагрузок (транзакции загрузки) и хранилищ (транзакции хранилища).Это 32-байтовые транзакции.В любом случае, загружает или хранит, транзакции являются и должны быть полностью эффективными в этом случае.Запрашиваемый трафик соответствует фактическому трафику, и другие метрики профилировщика подтверждают это.Если фактический трафик был намного выше, чем запрошенный трафик, это было бы хорошим показателем неэффективных нагрузок или хранилищ:

  1                  gld_requested_throughput                                      Requested Global Load Throughput  150.32GB/s  150.32GB/s  150.32GB/s
  1                  gst_requested_throughput                                     Requested Global Store Throughput  150.32GB/s  150.32GB/s  150.32GB/s
  1                            gld_throughput                                                Global Load Throughput  150.32GB/s  150.32GB/s  150.32GB/s
  1                            gst_throughput                                               Global Store Throughput  150.32GB/s  150.32GB/s  150.32GB/s

Запрошенная и достигнутая глобальная пропускная способность загрузки / хранения (gld_requested_throughput, gst_requested_throughput, gld_throughput, gst_throughput) составляет 150,32 ГБ / с каждый.Я ожидал бы более низкую пропускную способность для нагрузок, чем для магазинов, так как у нас больше транзакций на запрос.

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

gld_transactions = 536870914, но l2_read_transactions = 134218800. Глобальная память всегдадоступ через кэш L1 / L2.Почему количество транзакций чтения L2 намного меньше?Это не может быть все кэшировано в L1.(global_hit_rate = 0%)

Это просто из-за разного размера транзакций. Вы уже подсчитали, что кажущийся размер транзакции глобальной загрузки составляет 8 байт, и я уже указывал, что размер транзакции L2 составляет 32 байта, поэтому имеет смысл соотношение 4: 1 между общим количеством транзакций. , поскольку они отражают одно и то же движение одних и тех же данных при просмотре через 2 разных объектива. Обратите внимание, что всегда было несоответствие между размером глобальных транзакций и размером транзакций L2 или транзакций в DRAM. Просто их соотношение может варьироваться в зависимости от архитектуры графического процессора и, возможно, других факторов, таких как схемы загрузки.

Некоторые заметки:

  • Я не смогу ответить на такие вопросы, как «почему это так?» Или «почему Паскаль изменился с Ферми / Кеплера?» или «учитывая этот конкретный код, что бы вы прогнозировали как необходимые глобальные транзакции загрузки на этом конкретном графическом процессоре?», или «как правило, для этого конкретного графического процессора, как бы я рассчитывал или прогнозировал размер транзакции?»

  • Кроме того, NVIDIA предлагает новые инструменты профилирования (Nsight Compute и Nsight Systems) для работы с графическими процессорами. Многие из показателей эффективности и транзакций по запросу, которые доступны в nvprof , пропали в новой цепочке инструментов. Таким образом, эти установки в любом случае должны быть нарушены, потому что эти методы определения эффективности не будут доступны в будущем, основываясь на текущем наборе метрик.

  • Обратите внимание, что использование параметров компиляции, таких как -Xptxas -dlcm=ca, может повлиять на (L1) поведение кэширования. Однако я не ожидаю, что кэши окажут значительное влияние на производительность или эффективность для этого конкретного кода копии.

  • Это возможное уменьшение размера транзакции, как правило, хорошо. Это не приводит к потере эффективности для шаблонов трафика, таких как представленные в этом коде, а для некоторых других кодов позволяет (менее 128-байтных) запросов удовлетворяться при меньшей потере пропускной способности.

  • Хотя это не совсем Паскаль, здесь - это более точный пример возможной изменчивости этих измерений для Максвелла. Паскаль будет иметь аналогичную изменчивость. Кроме того, небольшая подсказка об этом изменении (особенно для Паскаля) была дана в Руководстве по настройке Паскаля . Он ни в коем случае не предлагает полного описания и не объясняет все ваши наблюдения, но намекает на общую идею о том, что глобальные транзакции больше не фиксируются на 128-байтовом размере.

...