*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
* * Neither the name of NVIDIA CORPORATION nor the names of its
* contributors may be used to endorse or promote products derived
* from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/
Bindless Texture/Surface
This sample generates a few 2D textures and uses cudaTextureObjects to
perform pseudo virtual texturing for display. One 2D texture stores
references to other textures.
Furthermore use of mip mapping is shown using both cudaTextureObjects
and cudaSurfaceObjects.
Look into the bindlessTexture_kernel.cu file for most relevant code.
*/
#include <helper_gl.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <vector>
#if defined(__APPLE__) || defined(MACOSX)
#pragma clang diagnostic ignored "-Wdeprecated-declarations"
#include <GLUT/glut.h>
#ifndef glutCloseFunc
#define glutCloseFunc glutWMCloseFunc
#endif
#else
#include <GL/freeglut.h>
#endif
#include <cuda_gl_interop.h>
#include <helper_functions.h>
#include "bindlessTexture.h"
#define MAX_EPSILON_ERROR 5.0f
#define THRESHOLD 0.15f
const char *sSDKsample = "CUDA bindlessTexture";
const char *imageFilenames[] = {
"flower.ppm",
"person.ppm",
"sponge.ppm",
};
const cudaExtent atlasSize = make_cudaExtent(4, 4, 0);
const dim3 windowSize(512, 512);
const dim3 windowBlockSize(16, 16, 1);
const dim3 windowGridSize(windowSize.x / windowBlockSize.x, windowSize.y / windowBlockSize.y);
float lod = 0.5;
GLuint pbo;
struct cudaGraphicsResource *cuda_pbo_resource = NULL;
bool animate = true;
StopWatchInterface *timer = NULL;
uint *d_output = NULL;
const int frameCheckNumber = 4;
int fpsCount = 0;
int fpsLimit = 1;
int g_Index = 0;
unsigned int frameCount = 0;
unsigned int g_TotalErrors = 0;
int *pArgc = NULL;
char **pArgv = NULL;
extern "C" void initAtlasAndImages(const Image *images, size_t numImages, cudaExtent atlasSize);
extern "C" void deinitAtlasAndImages();
extern "C" void randomizeAtlas();
extern "C" void renderAtlasImage(dim3 gridSize, dim3 blockSize, uint *d_output, uint imageW, uint imageH, float lod);
void computeFPS()
{
frameCount++;
fpsCount++;
if (fpsCount == fpsLimit) {
char fps[256];
float ifps = 1.f / (sdkGetAverageTimerValue(&timer) / 1000.f);
sprintf(fps, "%s: %3.1f fps", sSDKsample, ifps);
glutSetWindowTitle(fps);
fpsCount = 0;
fpsLimit = (int)MAX(1.0f, ifps);
sdkResetTimer(&timer);
}
}
void render()
{
checkCudaErrors(cudaGraphicsMapResources(1, &cuda_pbo_resource, 0));
size_t num_bytes;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&d_output, &num_bytes, cuda_pbo_resource));
renderAtlasImage(windowGridSize, windowBlockSize, d_output, windowSize.x, windowSize.y, lod);
getLastCudaError("render_kernel failed");
checkCudaErrors(cudaGraphicsUnmapResources(1, &cuda_pbo_resource, 0));
}
void display()
{
sdkStartTimer(&timer);
render();
glClear(GL_COLOR_BUFFER_BIT);
glDisable(GL_DEPTH_TEST);
glRasterPos2i(0, 0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
glDrawPixels(windowSize.x, windowSize.y, GL_RGBA, GL_UNSIGNED_BYTE, 0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
glutSwapBuffers();
glutReportErrors();
sdkStopTimer(&timer);
computeFPS();
}
void idle()
{
if (animate) {
lod += 0.02f;
glutPostRedisplay();
}
}
void keyboard(unsigned char key, int x, int y)
{
switch (key) {
case 27:
#if defined(__APPLE__) || defined(MACOSX)
exit(EXIT_SUCCESS);
#else
glutDestroyWindow(glutGetWindow());
return;
#endif
break;
case '=':
case '+':
lod += 0.25f;
break;
case '-':
lod -= 0.25f;
break;
case 'r':
randomizeAtlas();
break;
case ' ':
animate = !animate;
lod = 0.0f;
break;
default:
break;
}
glutPostRedisplay();
}
void reshape(int x, int y)
{
glViewport(0, 0, x, y);
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(0.0, 1.0, 0.0, 1.0, 0.0, 1.0);
}
void cleanup()
{
sdkDeleteTimer(&timer);
if (cuda_pbo_resource) {
cudaGraphicsUnregisterResource(cuda_pbo_resource);
glDeleteBuffers(1, &pbo);
}
}
void cleanup_all()
{
cleanup();
deinitAtlasAndImages();
}
void initGLBuffers()
{
glGenBuffers(1, &pbo);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, windowSize.x * windowSize.y * sizeof(GLubyte) * 4, 0, GL_STREAM_DRAW_ARB);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
checkCudaErrors(cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo, cudaGraphicsMapFlagsWriteDiscard));
}
uchar *loadRawFile(const char *filename, size_t size)
{
FILE *fp = fopen(filename, "rb");
if (!fp) {
fprintf(stderr, "Error opening file '%s'\n", filename);
return 0;
}
uchar *data = (uchar *)malloc(size);
size_t read = fread(data, 1, size, fp);
fclose(fp);
printf("Read '%s', %zu bytes\n", filename, read);
return data;
}
void initGL(int *argc, char **argv)
{
glutInit(argc, argv);
glutInitDisplayMode(GLUT_RGB | GLUT_DOUBLE);
glutInitWindowSize(windowSize.x, windowSize.y);
glutCreateWindow(sSDKsample);
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutReshapeFunc(reshape);
glutIdleFunc(idle);
if (!isGLVersionSupported(2, 0) || !areGLExtensionsSupported("GL_ARB_pixel_buffer_object")) {
fprintf(stderr, "Required OpenGL extensions are missing.");
exit(EXIT_FAILURE);
}
}
void runAutoTest(const char *ref_file, char *exec_path)
{
size_t windowBytes = windowSize.x * windowSize.y * sizeof(GLubyte) * 4;
checkCudaErrors(cudaMalloc((void **)&d_output, windowBytes));
renderAtlasImage(windowGridSize, windowBlockSize, d_output, windowSize.x, windowSize.y, lod);
checkCudaErrors(cudaDeviceSynchronize());
getLastCudaError("render_kernel failed");
void *h_output = malloc(windowBytes);
checkCudaErrors(cudaMemcpy(h_output, d_output, windowBytes, cudaMemcpyDeviceToHost));
sdkDumpBin(h_output, (unsigned int)windowBytes, "bindlessTexture.bin");
bool bTestResult = sdkCompareBin2BinFloat("bindlessTexture.bin",
sdkFindFilePath(ref_file, exec_path),
windowSize.x * windowSize.y,
MAX_EPSILON_ERROR,
THRESHOLD,
exec_path);
checkCudaErrors(cudaFree(d_output));
free(h_output);
deinitAtlasAndImages();
sdkStopTimer(&timer);
sdkDeleteTimer(&timer);
exit(bTestResult ? EXIT_SUCCESS : EXIT_FAILURE);
}
void loadImageData(const char *exe_path)
{
std::vector<Image> images;
for (size_t i = 0; i < sizeof(imageFilenames) / sizeof(imageFilenames[0]); i++) {
unsigned int imgWidth = 0;
unsigned int imgHeight = 0;
uchar *imgData = NULL;
const char *imgPath = 0;
const char *imgFilename = imageFilenames[i];
if (exe_path) {
imgPath = sdkFindFilePath(imgFilename, exe_path);
}
if (imgPath == 0) {
printf("Error finding image file '%s'\n", imgFilename);
exit(EXIT_FAILURE);
}
sdkLoadPPM4(imgPath, (unsigned char **)&imgData, &imgWidth, &imgHeight);
if (!imgData) {
printf("Error opening file '%s'\n", imgPath);
exit(EXIT_FAILURE);
}
printf("Loaded '%s', %d x %d pixels\n", imgPath, imgWidth, imgHeight);
checkHost(imgWidth > 1);
checkHost(imgHeight > 1);
Image img;
img.size = make_cudaExtent(imgWidth, imgHeight, 0);
img.h_data = imgData;
images.push_back(img);
}
initAtlasAndImages(&images[0], images.size(), atlasSize);
}
int main(int argc, char **argv)
{
sdkCreateTimer(&timer);
pArgc = &argc;
pArgv = argv;
char *ref_file = NULL;
#if defined(__linux__)
setenv("DISPLAY", ":0", 0);
#endif
printf("%s Starting...\n\n", sSDKsample);
if (checkCmdLineFlag(argc, (const char **)argv, "file")) {
fpsLimit = frameCheckNumber;
getCmdLineArgumentString(argc, (const char **)argv, "file", &ref_file);
}
srand(15234);
findCudaDevice(argc, (const char **)argv);
if (!ref_file) {
initGL(&argc, argv);
initGLBuffers();
}
if (!checkCudaCapabilities(3, 0)) {
cleanup();
exit(EXIT_WAIVED);
}
loadImageData(argv[0]);
if (ref_file) {
runAutoTest(ref_file, argv[0]);
}
printf("Press space to toggle animation\n"
"Press '+' and '-' to change lod level\n"
"Press 'r' to randomize virtual atlas\n");
#if defined(__APPLE__) || defined(MACOSX)
atexit(cleanup_all);
#else
glutCloseFunc(cleanup_all);
#endif
glutMainLoop();
}