mindfusion的博客

include_directories(/usr/local/cuda/include)link_directories(/usr/local/cuda/lib64)

nvprof --metrics achieved_occupancy,gld_throughput,gst_throughput,gld_efficiency,gst_efficiency,gld_transactions,gst_transactions,gld_transactions_per_request,gst_transactions_per_request,branch_effic...

CUDA SHARED MEMORY在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。GPU上的memory有两种：· On-board memory· On-chip memoryglo...

一般使用__ldg是更好的选择。通过constant缓存存储的数据必须相对较小而且必须获取同一个地址以便获取最佳性能，相反，只读缓存则可以存放较大的数据，且不必地址一致。__global__ void transpose2(const real* A, real* B, const int N){ const int nx = blockIdx.x * blockDim.x + ...

__ ldg（） 内置仅适用于计算能力3.5（或更高版本）架构。这意味着：必须在计算3.5（或更新版本）GPU上运行 必须针对计算3.5（或更新版本）的GPU进行编译 不能为旧体系结构编译。解决方式修改cmake:set(CUDA_ARCH "-arch=sm_50" CACHE STRING "Value of the NVCC -arch option."...

CUDA Memory Model对于程序员来说，memory可以分为下面两类：Programmable：我们可以灵活操作的部分。 Non-programmable：不能操作，由一套自动机制来达到很好的性能。在CPU的存储结构中，L1和L2 cache都是non-programmable的。对于CUDA来说，programmable的类型很丰富：Registers Shared ...

mian.cu#include <stdio.h>#define CHECK(call) \do \{ ...

学到了cudaMemcpyToSymbol竟然还有将数据从host拷贝到global的功能，以前只用过这个函数拷贝到constant memory。拷贝方式的不同是由目的内存申请的方式决定的。申请的是device内存，cudaMemcpyToSymbol拷贝就是从host拷贝到global memory。申请的是constant内存，cudaMemcpyToSymbol拷贝就是从ho...

https://bbs.gpuworld.cn/

添加头文件：#include <intrin.h>

参考书籍：《cuda专家手册|GPU编程权威》1：_shfl_xor首先介绍__shfl_xor，因为最先用到它。__shfl_xor（var,laneMask）：Copy from a lane based on bitwise XOR of own lane ID意思就是从当前的线程id与laneMak异或运算的值作为线程号的，把这个线程号的var值取出来。演示图：...

如果cuda需要调试需要在配置cmake:cmake配置set(CUDA_NVCC_FLAGS-G;-g)命令行编译配置：nvcc –g –GMatrixCUDA.cu –oMatrixCUDA打开 nvidia nsight monitor ->options进行设置2.尽量把WDDM TDR Delay 调大点3.将 use this monit...

ubuntu环境上，安装cuda,会自动安装一些工具nvcc ,nvprof nvcc-gdb....1.编译生成可执行文件2.nvprof ./test可以把性能分析数据输出到文件中nvprof ./test -o test.nvvp可以把在nvidia visual profiler现实更加直观一些常用的配置参数： achieved_occupancy参...

cudasift