Replace manual addressing with cuda texture

Author: tigert1998

ray-tracing-cuda/model.h

@@ -82,10 +82,10 @@ struct Model {
                                  "/textures/" + BaseName(path.C_Str());
         LOG(INFO) << "loading texture at: \"" << image_path << "\"";
         auto data = stbi_load(image_path.c_str(), &textures[texture_id].width,
-                              &textures[texture_id].height, &channels, 3);
+                              &textures[texture_id].height, &channels, 4);
         textures[texture_id].data =
             std::string(data, data + (textures[texture_id].width *
-                                      textures[texture_id].height * 3));
+                                      textures[texture_id].height * 4));
         stbi_image_free(data);
         meshes[id].texture_id = texture_id;
       }

ray-tracing-cuda/textures/image_texture.cu

@@ -2,23 +2,37 @@
 
 using glm::clamp;
 
-__host__ __device__ ImageTexture::ImageTexture(int height, int width,
-                                               int components, uint8_t *data)
-    : height_(height), width_(width), components_(components), data_(data) {}
+__host__ __device__
+ImageTexture::ImageTexture(cudaTextureObject_t image_texture)
+    : image_texture_(image_texture) {}
 
 __device__ glm::vec3 ImageTexture::Value(double u, double v,
                                          const glm::vec3 &p) const {
-  u = clamp(u, 0.0, 1.0);
-  v = 1.0 - clamp(v, 0.0, 1.0);  // Flip V to image coordinates
+  v = 1.0 - v;  // Flip V to image coordinates
 
-  auto i = static_cast<int>(u * width_);
-  auto j = static_cast<int>(v * height_);
-
-  // Clamp integer mapping, since actual coordinates should be less than 1.0
-  if (i >= width_) i = width_ - 1;
-  if (j >= height_) j = height_ - 1;
-
-  uint8_t *pixel = data_ + j * (components_ * width_) + i * components_;
+  float4 ret = tex2D<float4>(image_texture_, u, v);
+  return glm::vec3(ret.x, ret.y, ret.z);
+}
 
-  return glm::vec3(pixel[0], pixel[1], pixel[2]) / 255.0f;
+cudaTextureObject_t ImageTexture::CreateCudaTextureObj(
+    uint8_t *dev_buffer, int height, int width, uint64_t pitch_in_bytes) {
+  cudaResourceDesc res_desc;
+  memset(&res_desc, 0, sizeof(res_desc));
+  res_desc.resType = cudaResourceTypePitch2D;
+  res_desc.res.pitch2D.devPtr = dev_buffer;
+  res_desc.res.pitch2D.desc.f = cudaChannelFormatKindUnsigned;
+  res_desc.res.pitch2D.desc.x = 8;
+  res_desc.res.pitch2D.desc.y = 8;
+  res_desc.res.pitch2D.desc.z = 8;
+  res_desc.res.pitch2D.desc.w = 8;
+  res_desc.res.pitch2D.height = height;
+  res_desc.res.pitch2D.width = width;
+  res_desc.res.pitch2D.pitchInBytes = pitch_in_bytes;
+  cudaTextureDesc tex_desc;
+  memset(&tex_desc, 0, sizeof(tex_desc));
+  tex_desc.readMode = cudaReadModeNormalizedFloat;
+  tex_desc.normalizedCoords = 1;
+  cudaTextureObject_t tex;
+  cudaCreateTextureObject(&tex, &res_desc, &tex_desc, nullptr);
+  return tex;
 }

ray-tracing-cuda/textures/image_texture.cuh

@@ -1,16 +1,20 @@
 #pragma once
 
+#include <cuda_runtime.h>
+
 #include <glm/glm.hpp>
 
 #include "textures/texture.cuh"
 
 class ImageTexture : public Texture {
  private:
-  uint8_t *data_;
-  int height_, width_, components_;
+  cudaTextureObject_t image_texture_;
 
  public:
-  __host__ __device__ ImageTexture(int height, int width, int components,
-                                   uint8_t *data);
+  __host__ __device__ ImageTexture(cudaTextureObject_t image_texture);
   __device__ glm::vec3 Value(double u, double v, const glm::vec3 &p) const;
+
+  static cudaTextureObject_t CreateCudaTextureObj(uint8_t *dev_buffer,
+                                                  int height, int width,
+                                                  uint64_t pitch_in_bytes);
 };

scenes/birthday.cu

@@ -31,21 +31,16 @@ curandState *d_states;
 Camera *d_camera;
 HitableList *d_world;
 glm::vec3 *d_image;
-
-struct DeviceImage {
-  uint8_t *image;
-  int height, width, components;
-};
-
-DeviceImage d_earthmap;
+uint8_t *d_tex_image;
+cudaTextureObject_t tex;
 
 using glm::pi;
 using glm::rotateX;
 using glm::rotateY;
 using glm::vec3;
 
 __global__ void InitWorld(HitableList *world, Camera *camera,
-                          DeviceImage earthmap) {
+                          cudaTextureObject_t tex) {
   new (world) HitableList();
   new (camera) Camera(vec3(278, 278, -800), vec3(278, 278, 0), vec3(0, 1, 0),
                       pi<double>() * 2 / 9, double(WIDTH) / HEIGHT);
@@ -56,8 +51,7 @@ __global__ void InitWorld(HitableList *world, Camera *camera,
 
   auto light_material_ptr =
       new DiffuseLight(new ConstantTexture(vec3(1, 1, 1)));
-  auto earthmap_texture_ptr = new ImageTexture(
-      earthmap.height, earthmap.width, earthmap.components, earthmap.image);
+  auto earthmap_texture_ptr = new ImageTexture(tex);
   auto earthmap_material_ptr = new Lambertian(earthmap_texture_ptr);
   auto sky = new Sky();
 
@@ -79,22 +73,27 @@ __global__ void InitWorld(HitableList *world, Camera *camera,
   world->Append(new Sphere(vec3(278, 278, 0), 100, earthmap_material_ptr));
 }
 
-void InitImageTextures(DeviceImage *earthmap) {
-  auto data = stbi_load("resources/earthmap.jpg", &earthmap->width,
-                        &earthmap->height, &earthmap->components, 0);
-  int size = earthmap->height * earthmap->width * earthmap->components;
-  cudaMalloc(&earthmap->image, size);
-  cudaMemcpy(earthmap->image, data, size, cudaMemcpyHostToDevice);
+cudaTextureObject_t InitImageTextures() {
+  int width, height, channels;
+  uint64_t pitch_in_bytes;
+  auto data =
+      stbi_load("resources/earthmap.jpg", &width, &height, &channels, 4);
+  cudaMallocPitch(&d_tex_image, &pitch_in_bytes, 4 * width, height);
+  cudaMemcpy2D(d_tex_image, pitch_in_bytes, data, 4 * width, 4 * width, height,
+               cudaMemcpyHostToDevice);
+  auto tex = ImageTexture::CreateCudaTextureObj(d_tex_image, height, width,
+                                                pitch_in_bytes);
   auto err = cudaGetLastError();
   CHECK(err == cudaSuccess) << cudaGetErrorString(err);
+  return tex;
 }
 
 int main() {
   Main(
       &d_states, &d_camera, &d_world, &d_image,
       [](HitableList *world, Camera *camera) {
-        InitImageTextures(&d_earthmap);
-        InitWorld<<<1, 1>>>(world, camera, d_earthmap);
+        auto tex = InitImageTextures();
+        InitWorld<<<1, 1>>>(world, camera, tex);
       },
       HEIGHT, WIDTH, 200);
   return 0;