Накладные расходы Cuda L2

Question

У меня есть ядро ​​для проверки точек рендеринга с atomicMin. У тестовой установки есть тонны точек в макете памяти случая. Два буфера, один uint32 для кластеров 256x uint32.

namespace Point 
{ 
struct PackedBitfield 
{ 
    glm::uint32_t x : 6; 
    glm::uint32_t y : 6; 
    glm::uint32_t z : 6; 
    glm::uint32_t nx : 4; 
    glm::uint32_t ny : 4; 
    glm::uint32_t nz : 4; 
    glm::uint32_t unused : 2; 
}; 

union __align__(4) Packed 
{ 
    glm::uint32_t bits; 
    PackedBitfield field; 
}; 

struct ClusterPositionBitfield 
{ 
    glm::uint32_t x : 10; 
    glm::uint32_t y : 10; 
    glm::uint32_t z : 10; 
    glm::uint32_t w : 2; 
}; 

union ClusterPosition 
{ 
    glm::uint32_t bits; 
    ClusterPositionBitfield field; 
}; 
} 

// 
// launch with blockSize=(256, 1, 1) and grid=(numberOfClusters, 1, 1) 
// 
extern "C" __global__ void pointsRenderKernel(mat4 u_mvp, 
        ivec2 u_resolution, 
        uint64_t* rasterBuffer, 
        Point::Packed* points, 
        Point::ClusterPosition* clusterPosition) 
{ 
// extract and compute world position 
const Point::ClusterPosition cPosition(clusterPosition[blockIdx.x]); 
const Point::Packed point(points[blockIdx.x*256 + threadIdx.x]); 

...use points and write to buffer... 

} 

В результате SASS выглядит следующим образом:

Посмотрите на выходе профилировщика памяти: передача служебной L2 из Point::Packed* буфера чтения является 3,0. Почему? Память должна быть идеально выровнена и последовательна. Также почему это автоматически генерирует LDG (compute_50, sm_50)? Мне это не нужно кэшировать.

Answer 1

В подсказке для L2 передачи Накладные, он говорит, что он измеряет «количество байтов, передаваемых между L1 и L2 для каждого запрашиваемого байта в L1», и он также говорит, что «ниже, тем лучше».

В моем случае L2 Transfer Uphead для чтения Point::Packed - 1.0.

TEST CODE

namespace Point 
{ 
    struct PackedBitfield 
    { 
     uint32_t x : 6; 
     uint32_t y : 6; 
     uint32_t z : 6; 
     uint32_t nx : 4; 
     uint32_t ny : 4; 
     uint32_t nz : 4; 
     uint32_t unused : 2; 
    }; 

    union __align__(4) Packed 
    { 
     uint32_t bits; 
     PackedBitfield field; 
    }; 

    struct ClusterPositionBitfield 
    { 
     uint32_t x : 10; 
     uint32_t y : 10; 
     uint32_t z : 10; 
     uint32_t w : 2; 
    }; 

    union ClusterPosition 
    { 
     uint32_t bits; 
     ClusterPositionBitfield field; 
    }; 
} 

__global__ void pointsRenderKernel(Point::Packed* points, Point::ClusterPosition* clusterPosition) 
{ 
    int t_id = blockIdx.x * blockDim.x + threadIdx.x; 

    clusterPosition[blockIdx.x + blockDim.x] = clusterPosition[blockIdx.x]; 
    points[t_id + blockDim.x * gridDim.x] = points[t_id]; 
} 

void main() 
{ 
    int blockSize = 256; 
    int numberOfClusters = 256; 

    std::cout << sizeof(Point::Packed) << std::endl; 
    std::cout << sizeof(Point::ClusterPosition) << std::endl; 

    Point::Packed *d_points; 
    cudaMalloc(&d_points, sizeof(Point::Packed) * numberOfClusters * blockSize * 2); 

    Point::ClusterPosition *d_clusterPositions; 
    cudaMalloc(&d_points, sizeof(Point::ClusterPosition) * numberOfClusters * 2); 

    pointsRenderKernel<<<numberOfClusters, blockSize>>>(d_points, d_clusterPositions); 
} 

---

UPDATE

У меня были другие проблемы с Nsight раньше, когда я использовал последнюю версию драйвера. Я понизил драйвер до версии, поставляемой с установщиком CUDA 8.0.61 по умолчанию (скачан с here), и это устранило проблему. Версия, поставляемая с установщиком, - 376,51. Протестировано на Windows 10 64-bit и Visual Studio 2015, версия Nsight - 5,2, а моя графическая карта - cc6.1.

Это моя полная команда компилятора:

nvcc.exe -gencode = арка = compute_61, код = \ "sm_61, compute_61 \" --use-местном-окр --cl-версия 2015 - Xcompiler "/ wd 4819" -ccbin "C: \ Program Files (x86) \ Microsoft Visual Studio 14.0 \ VC \ bin \ x86_amd64" -I "C: \ Program Files \ NVIDIA GPU Computing Toolkit \ CUDA \ v8.0 \ include "-lineinfo --keep-dir x64 \ Release -maxrregcount = 0 --machine 64 --compile -cudart static -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler"/EHsc/W3/nologo/O2/FS/Zi/MD "-o x64 \ Release \ kernel.cu.obj kernel.cu"

---

UPDATE 2

я получаю тот же результат при компиляции с опцией sm_50,compute_50: 1.0 для L2 передачи служебной информации.