c++ - usar - ¿Cómo se mide el ancho de banda máximo de memoria en OpenGL?
uso del osciloscopio digital (1)
Le está pidiendo al controlador que lea de su memoria de proceso, dat . Esto provoca un tráfico de coherencia de caché extenso. Cuando la GPU lee esa memoria, no puede estar seguro de que esté actualizada, podría estar en la memoria caché de la CPU, modificada, y aún no escrita en la RAM. Esto hace que la GPU tenga que leer realmente desde la memoria caché de la CPU, que es mucho más costosa que pasar por alto la CPU y leer la memoria RAM. La RAM suele estar inactiva durante el funcionamiento normal, porque la tasa de aciertos de una CPU moderna suele ser del 95% al 99%. El caché se utiliza continuamente.
Para lograr el máximo rendimiento, debe permitir que el controlador asigne la memoria. La memoria normal que utiliza su programa, como las variables globales y el montón, se asignan en la memoria de reescritura . La memoria asignada por el controlador generalmente se asignará como combinación de escritura o incachable , lo que elimina el tráfico de coherencia.
Los números máximos de ancho de banda anunciados se lograrán solo sin la sobrecarga de coherencia de caché.
Para permitir que el controlador lo asigne, use glBufferData con un nullptr para los datos.
Sin embargo, no todo es color de rosa, si logras obligar al conductor a usar un búfer de combinación de escritura de memoria del sistema. Las lecturas de la CPU a tales direcciones serán muy lentas. Las escrituras secuenciales están optimizadas por la CPU, pero las escrituras aleatorias causarán que la escritura del búfer de combinación se vacíe con frecuencia, lo que perjudicará el rendimiento.
Solo para tener una idea de qué tipo de velocidades debería esperar, he estado tratando de comparar la transferencia entre la memoria global y los sombreadores, en lugar de confiar en las hojas de especificaciones de la GPU. Sin embargo no puedo acercarme al máximo teórico. De hecho estoy fuera por un factor de 50 !
Estoy usando un GTX Titan X, que se dice que tiene 336.5GB / s . Controlador Linux x64 352.21.
Encontré un punto de referencia CUDA here que me da ~ 240–250GB / s (esto es más de lo que esperaba).
Estoy tratando de hacer coincidir exactamente lo que hacen con los shaders. He intentado shaders de vértices, shaders de cómputo, acceso a objetos de búfer a través de image_load_store y NV_shader_buffer_store , con float s, vec4 s, bucles dentro del shader (con direccionamiento unido dentro del grupo de trabajo) y varios métodos de sincronización. Estoy atascado en ~ 7GB / s ( ver la actualización a continuación ).
¿Por qué GL es mucho más lento? ¿Estoy haciendo algo mal y, de ser así, cómo debería hacerse?
Aquí está mi MWE con tres métodos (1. sombreado de vértice con image_load_store, 2. sombreador de vértice con gráficos sin encuadernación, 3. sombreador de cálculo con gráficos sin encuadernación):
//#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 : enable/n"
"uniform float* data;/n"
"uniform float val;/n"
"void main() {/n"
" data[gl_VertexID] = val;/n"
" gl_Position = vec4(0.0);/n"
"}/n";
const char* bindlessComputeSource =
"#version 440/n"
"#extension GL_NV_gpu_shader5 : enable/n"
"#extension GL_NV_shader_buffer_load : enable/n"
"layout(local_size_x = 256) in;/n"
"uniform float* data;/n"
"uniform float val;/n"
"void main() {/n"
" data[gl_GlobalInvocationID.x] = val;/n"
"}/n";
GLuint compile(GLenum type, const char* shaderSrc)
{
GLuint shader = glCreateShader(type);
glShaderSource(shader, 1, (const GLchar**)&shaderSrc, NULL);
glCompileShader(shader);
int success = 0;
int loglen = 0;
glGetShaderiv(shader, GL_COMPILE_STATUS, &success);
glGetShaderiv(shader, GL_INFO_LOG_LENGTH, &loglen);
GLchar* log = new GLchar[loglen];
glGetShaderInfoLog(shader, loglen, &loglen, log);
if (!success)
{
printf("%s/n", log);
exit(0);
}
GLuint program = glCreateProgram();
glAttachShader(program, shader);
glLinkProgram(program);
return program;
}
GLuint timerQueries[2];
void start()
{
glGenQueries(2, timerQueries);
glQueryCounter(timerQueries[0], GL_TIMESTAMP);
}
float stop()
{
glMemoryBarrier(GL_ALL_BARRIER_BITS);
GLsync sync = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
glWaitSync(sync, 0, GL_TIMEOUT_IGNORED);
glQueryCounter(timerQueries[1], GL_TIMESTAMP);
GLint available = 0;
while (!available) //sometimes gets stuck here for whatever reason
glGetQueryObjectiv(timerQueries[1], GL_QUERY_RESULT_AVAILABLE, &available);
GLuint64 a, b;
glGetQueryObjectui64v(timerQueries[0], GL_QUERY_RESULT, &a);
glGetQueryObjectui64v(timerQueries[1], GL_QUERY_RESULT, &b);
glDeleteQueries(2, timerQueries);
return b - a;
}
int main(int argc, char** argv)
{
float* check;
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB | GLUT_DEPTH);
glutCreateWindow("test");
glewInit();
int bufferSize = 64 * 1024 * 1024; //64MB
int loops = 500;
glEnable(GL_RASTERIZER_DISCARD);
float* dat = new float[bufferSize/sizeof(float)];
memset(dat, 0, bufferSize);
//create a buffer with data
GLuint buffer;
glGenBuffers(1, &buffer);
glBindBuffer(GL_TEXTURE_BUFFER, buffer);
glBufferData(GL_TEXTURE_BUFFER, bufferSize, NULL, GL_STATIC_DRAW);
//get a bindless address
GLuint64 address;
glMakeBufferResidentNV(GL_TEXTURE_BUFFER, GL_READ_WRITE);
glGetBufferParameterui64vNV(GL_TEXTURE_BUFFER, GL_BUFFER_GPU_ADDRESS_NV, &address);
//make a texture alias for it
GLuint bufferTexture;
glGenTextures(1, &bufferTexture);
glBindTexture(GL_TEXTURE_BUFFER, bufferTexture);
glTexBuffer(GL_TEXTURE_BUFFER, GL_R32F, buffer);
glBindImageTextureEXT(0, bufferTexture, 0, GL_FALSE, 0, GL_READ_WRITE, GL_R32F);
//compile the shaders
GLuint imageShader = compile(GL_VERTEX_SHADER, imageSource);
GLuint bindlessShader = compile(GL_VERTEX_SHADER, bindlessSource);
GLuint bindlessComputeShader = compile(GL_COMPUTE_SHADER, bindlessComputeSource);
//warm-up and check values
glBufferData(GL_TEXTURE_BUFFER, bufferSize, dat, GL_STATIC_DRAW);
glUseProgram(imageShader);
glUniform1i(glGetUniformLocation(imageShader, "data"), 0);
glUniform1f(glGetUniformLocation(imageShader, "val"), 1.0f);
glDrawArrays(GL_POINTS, 0, bufferSize/sizeof(float));
glMemoryBarrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT);
//check = (float*)glMapBuffer(GL_TEXTURE_BUFFER, GL_READ_ONLY);
//for (int i = 0; i < bufferSize/sizeof(float); ++i)
// assert(check[i] == 1.0f);
//glUnmapBuffer(GL_TEXTURE_BUFFER);
glBufferData(GL_TEXTURE_BUFFER, bufferSize, dat, GL_STATIC_DRAW);
glUseProgram(bindlessShader);
glProgramUniformui64NV(bindlessShader, glGetUniformLocation(bindlessShader, "data"), address);
glUniform1f(glGetUniformLocation(bindlessShader, "val"), 1.0f);
glDrawArrays(GL_POINTS, 0, bufferSize/sizeof(float));
//glMemoryBarrier(GL_ALL_BARRIER_BITS); //this causes glDispatchCompute to segfault later, so don''t uncomment
//check = (float*)glMapBuffer(GL_TEXTURE_BUFFER, GL_READ_ONLY);
//for (int i = 0; i < bufferSize/sizeof(float); ++i)
// assert(check[i] == 1.0f);
//glUnmapBuffer(GL_TEXTURE_BUFFER);
glBufferData(GL_TEXTURE_BUFFER, bufferSize, dat, GL_STATIC_DRAW);
glUseProgram(bindlessComputeShader);
glProgramUniformui64NV(bindlessComputeShader, glGetUniformLocation(bindlessComputeShader, "data"), address);
glUniform1f(glGetUniformLocation(bindlessComputeShader, "val"), 1.0f);
glDispatchCompute(bufferSize/(sizeof(float) * 256), 1, 1);
glMemoryBarrier(GL_ALL_BARRIER_BITS);
//check = (float*)glMapBuffer(GL_TEXTURE_BUFFER, GL_READ_ONLY);
//for (int i = 0; i < bufferSize/sizeof(float); ++i)
// assert(check[i] == 1.0f); //glDispatchCompute doesn''t actually write anything with bindless graphics
//glUnmapBuffer(GL_TEXTURE_BUFFER);
glFinish();
//time image_load_store
glUseProgram(imageShader);
glUniform1i(glGetUniformLocation(imageShader, "data"), 0);
glUniform1f(glGetUniformLocation(imageShader, "val"), 1.0f);
start();
for (int i = 0; i < loops; ++i)
glDrawArrays(GL_POINTS, 0, bufferSize/sizeof(float));
GLuint64 imageTime = stop();
printf("image_load_store: %.2fGB/s/n", (float)((bufferSize * (double)loops) / imageTime));
//time bindless
glUseProgram(bindlessShader);
glProgramUniformui64NV(bindlessShader, glGetUniformLocation(bindlessShader, "data"), address);
glUniform1f(glGetUniformLocation(bindlessShader, "val"), 1.0f);
start();
for (int i = 0; i < loops; ++i)
glDrawArrays(GL_POINTS, 0, bufferSize/sizeof(float));
GLuint64 bindlessTime = stop();
printf("bindless: %.2fGB/s/n", (float)((bufferSize * (double)loops) / bindlessTime));
//time bindless in a compute shader
glUseProgram(bindlessComputeShader);
glProgramUniformui64NV(bindlessComputeShader, glGetUniformLocation(bindlessComputeShader, "data"), address);
glUniform1f(glGetUniformLocation(bindlessComputeShader, "val"), 1.0f);
start();
for (int i = 0; i < loops; ++i)
glDispatchCompute(bufferSize/(sizeof(float) * 256), 1, 1);
GLuint64 bindlessComputeTime = stop();
printf("bindless compute: %.2fGB/s/n", (float)((bufferSize * (double)loops) / bindlessComputeTime));
assert(glGetError() == GL_NO_ERROR);
return 0;
}
Mi salida:
image_load_store: 6.66GB/s
bindless: 6.68GB/s
bindless compute: 6.65GB/s
Algunas notas:
- Los sombreadores informáticos con gráficos sin encuadernación no parecen escribir nada (la afirmación comentada falla), o al menos los datos no se recuperan con
glMapBufferaunque la velocidad coincida con los otros métodos. El uso de image_load_store en el sombreador de cómputo funciona y da la misma velocidad que los sombreadores de vértice (aunque pensé que sería una permutación demasiada para publicar). - Llamar a
glMemoryBarrier(GL_ALL_BARRIER_BITS)antes de queglDispatchComputecause un bloqueo en el controlador. - Comentando los tres
glBufferData(GL_TEXTURE_BUFFER, bufferSize, dat, GL_STATIC_DRAW);, que se usa para verificar la salida, aumenta la velocidad de las dos primeras pruebas a 17GB / sy el shader de cómputo se dispara a 292GB / s, lo cual está mucho más cerca de lo que me gustaría, pero esto no se puede confiar debido al punto 1. - A veces,
while (!available)cuelga durante años (ctrl-c cuando me canso de esperar, muestra que todavía está en el circuito).
Para referencia, aquí está el código CUDA:
//http://www.ks.uiuc.edu/Research/vmd/doxygen/CUDABench_8cu-source.html
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cuda.h>
#define CUERR { cudaError_t err; /
if ((err = cudaGetLastError()) != cudaSuccess) { /
printf("CUDA error: %s, %s line %d/n", cudaGetErrorString(err), __FILE__, __LINE__); /
return -1; }}
//
// GPU device global memory bandwidth benchmark
//
template <class T>
__global__ void gpuglobmemcpybw(T *dest, const T *src) {
const unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x;
dest[idx] = src[idx];
}
template <class T>
__global__ void gpuglobmemsetbw(T *dest, const T val) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
dest[idx] = val;
}
typedef float4 datatype;
static int cudaglobmembw(int cudadev, double *gpumemsetgbsec, double *gpumemcpygbsec) {
int i;
int len = 1 << 22; // one thread per data element
int loops = 500;
datatype *src, *dest;
datatype val=make_float4(1.0f, 1.0f, 1.0f, 1.0f);
// initialize to zero for starters
float memsettime = 0.0f;
float memcpytime = 0.0f;
*gpumemsetgbsec = 0.0;
*gpumemcpygbsec = 0.0;
// attach to the selected device
cudaError_t rc;
rc = cudaSetDevice(cudadev);
if (rc != cudaSuccess) {
#if CUDART_VERSION >= 2010
rc = cudaGetLastError(); // query last error and reset error state
if (rc != cudaErrorSetOnActiveProcess)
return -1; // abort and return an error
#else
cudaGetLastError(); // just ignore and reset error state, since older CUDA
// revs don''t have a cudaErrorSetOnActiveProcess enum
#endif
}
cudaMalloc((void **) &src, sizeof(datatype)*len);
CUERR
cudaMalloc((void **) &dest, sizeof(datatype)*len);
CUERR
dim3 BSz(256, 1, 1);
dim3 GSz(len / (BSz.x * BSz.y * BSz.z), 1, 1);
// do a warm-up pass
gpuglobmemsetbw<datatype><<< GSz, BSz >>>(src, val);
CUERR
gpuglobmemsetbw<datatype><<< GSz, BSz >>>(dest, val);
CUERR
gpuglobmemcpybw<datatype><<< GSz, BSz >>>(dest, src);
CUERR
cudaEvent_t start, end;
cudaEventCreate(&start);
cudaEventCreate(&end);
// execute the memset kernel
cudaEventRecord(start, 0);
for (i=0; i<loops; i++) {
gpuglobmemsetbw<datatype><<< GSz, BSz >>>(dest, val);
}
CUERR
cudaEventRecord(end, 0);
CUERR
cudaEventSynchronize(start);
CUERR
cudaEventSynchronize(end);
CUERR
cudaEventElapsedTime(&memsettime, start, end);
CUERR
// execute the memcpy kernel
cudaEventRecord(start, 0);
for (i=0; i<loops; i++) {
gpuglobmemcpybw<datatype><<< GSz, BSz >>>(dest, src);
}
cudaEventRecord(end, 0);
CUERR
cudaEventSynchronize(start);
CUERR
cudaEventSynchronize(end);
CUERR
cudaEventElapsedTime(&memcpytime, start, end);
CUERR
cudaEventDestroy(start);
CUERR
cudaEventDestroy(end);
CUERR
*gpumemsetgbsec = (len * sizeof(datatype) / (1024.0 * 1024.0)) / (memsettime / loops);
*gpumemcpygbsec = (2 * len * sizeof(datatype) / (1024.0 * 1024.0)) / (memcpytime / loops);
cudaFree(dest);
cudaFree(src);
CUERR
return 0;
}
int main()
{
double a, b;
cudaglobmembw(0, &a, &b);
printf("%f %f/n", (float)a, (float)b);
return 0;
}
Actualizar:
Parece que el búfer se hace no residente en mis llamadas glBufferData que estaban allí para comprobar que la salida se estaba escribiendo. Según la extensión :
Un búfer también se convierte en no residente de forma implícita como resultado de ser re-especificado a través de BufferData o de ser eliminado.
...
BufferData se especifica para "eliminar el almacén de datos existente", por lo que la dirección de la GPU de esos datos debería ser inválida. Por lo tanto, el búfer se hace no residente en el contexto actual.
En una conjetura, OpenGL luego transmite los datos del objeto del búfer en cada fotograma y no los almacena en la memoria de video. Esto explica por qué el shader de cálculo falló en la afirmación, sin embargo, existe una ligera anomalía en el hecho de que los gráficos sin enlace en el shader de vértice aún funcionaban cuando no eran residentes, pero por ahora lo ignoraré. No tengo idea de por qué un objeto de búfer de 64 MB no sería el residente (aunque quizás después del primer uso) cuando hay 12 GB disponibles.
Entonces, después de cada llamada a glBufferData lo hago residente de nuevo y obtengo la dirección en caso de que cambie:
glBufferData(GL_TEXTURE_BUFFER, bufferSize, dat, GL_STATIC_DRAW);
glMakeBufferResidentNV(GL_TEXTURE_BUFFER, GL_READ_WRITE);
glGetBufferParameterui64vNV(GL_TEXTURE_BUFFER, GL_BUFFER_GPU_ADDRESS_NV, &address);
assert(glIsBufferResidentNV(GL_TEXTURE_BUFFER)); //sanity check
Ahora estoy obteniendo 270–290GB / s con el sombreador de cómputo usando image_load_store o gráficos sin encuadernación. Ahora mi pregunta incluye :
- Dado que el búfer parece ser residente para cada prueba y el sombreador de cálculo es bueno y rápido, ¿por qué las versiones de sombreado de vértices son tan lentas?
Sin la extensión de gráficos sin encuadernación, ¿cómo deberían los usuarios regulares de OpenGL poner datos en la memoria de video (en realidad poner y no sugerir de forma ociosa que al controlador le gustaría)?
Estoy bastante seguro de que me habría dado cuenta de este problema en situaciones del mundo real, y es este punto de referencia ideado el que se mueve lentamente, así que, ¿cómo podría engañar al conductor para que haga que un objeto de búfer resida? Ejecutar un sombreador de cálculo primero no cambia nada.