相关疑难解决方法(0)

我试图了解手写内核的每个CUDA线程的资源使用情况.

我把我的kernel.cu文件编译成了一个kernel.o文件nvcc -arch=sm_20 -ptxas-options=-v

我得到了以下输出

ptxas info    : Compiling entry function 'searchkernel(octree, int*, double, int, double*, double*, double*)' for 'sm_20'
ptxas info    : Function properties for searchkernel(octree, int*, double, int, double*, double*, double*)
    72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14]

看看上面的输出,这是正确的

• 每个CUDA线程使用46个寄存器？
• 没有寄存器溢出到本地内存？

我在理解输出方面也遇到了一些问题.

• 我的内核调用了很多c++filt函数.IS 72字节是__device__和__global__函数堆栈帧的内存总和？

• __device__和之间有什么区别0 byte spill stores …

只是想知道我应该期待什么样的速度,我一直试图在全局内存和着色器之间进行基准测试,而不是依赖于GPU规格表.但是我无法接近理论上的最大值.事实上,我出局了50倍!

我正在使用GTX Titan X,据说它有336.5GB/s.Linux x64驱动程序352.21.

我在这里找到了一个CUDA基准测试,它给了我~240-250GB/s(这更符合我的预期).

我正在尝试将它们与着色器完全匹配.我已经试过顶点着色器,计算着色,经由访问缓冲器对象image_load_store和NV_shader_buffer_store,具有floatS,vec4S,着色器内循环(与工作组内聚结的寻址)和定时的各种方法.我卡在~7GB/s(参见下面的更新).

为什么GL这么慢？我做错了什么,如果是的话,应该怎么做？

这是我的MWE有三种方法(1.顶点着色器与image_load_store,2.顶点着色器与无绑定图形,3.计算着色器与无绑定图形):

//#include <windows.h>
#include <assert.h>
#include <stdio.h>
#include <memory.h>
#include <GL/glew.h>
#include <GL/glut.h>

const char* imageSource =
    "#version 440\n"
    "uniform layout(r32f) imageBuffer data;\n"
    "uniform float val;\n"
    "void main() {\n"
    "   imageStore(data, gl_VertexID, vec4(val, 0.0, 0.0, 0.0));\n"
    "   gl_Position = vec4(0.0);\n"
    "}\n";

const char* bindlessSource =
    "#version 440\n"
    "#extension GL_NV_gpu_shader5 : enable\n"
    "#extension GL_NV_shader_buffer_load : …

我的内核中有很多未使用的寄存器.我想告诉CUDA使用一些寄存器来保存一些数据,而不是每次需要时都读取全局数据.(我无法使用共享内存.)

__global__ void simple(float *gData) {
float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

编译瓦特/:NVCC -arch sm_20 --ptxas选项= -v simple.cu,我也得到
0字节堆栈帧,0字节溢出存储,0字节溢出负载
使用2个寄存器,40个字节CMEM [0]

__global__ void simple(float *gData) {
register float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

注册声明什么都不做.
0字节堆栈帧,0字节溢出存储,0字节溢出加载
使用2个寄存器,40字节cmem [0]

__global__ void simple(float *gData) {
volatile float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
} …