I am extending Roger Allen's Ray Tracing in One Weekend in CUDA (and also here) by following Peter Shirley's Ray Tracing: The Next Week; specifically Image Texture Mapping.
The image I obtain from CUDA is not as sharp as the CPU version. Shirley's book uses double
and Allen restricts to float
but could that be the only reason?
The texture source for both is a 24 bit depth jpeg file.
In both cases 100x sampling per pixel used for antialiasing.
Is there anything to be done in the kernel to improve the image quality?
class imageTextureX :public textureX {
public:
__device__ imageTextureX(uint32_t width, uint32_t height, cudaTextureObject_t textureObject) :_width(width), _height(height), _textureObj(textureObject) {}
__forceinline__ __device__ vec3 value(float u, float v, const vec3& p) const override {
//Flip Y-Coordinate;
v = 1.0f - v;
//De-normalize coordinate(s)
uint32_t X = u * _width;
//Re-normalize
constexpr float pixelOffset{ 0.5f };
const float denominator = 1.0f / (float)(3 * _width);
u = (float)(3 * X + pixelOffset);
vec3 colorVal;
colorVal.e[0] = (float)tex2D<uint8_t>(_textureObj, u++ * denominator, v);
colorVal.e[1] = (float)tex2D<uint8_t>(_textureObj, u++ * denominator, v);
colorVal.e[2] = (float)tex2D<uint8_t>(_textureObj, u * denominator, v);
return colorVal / 255.0f;
}
private:
cudaTextureObject_t _textureObj;
uint32_t _width;
uint32_t _height;
};
Texture definition code
//Load Texture Image from Disk
cudaArray* cuArray;
cudaTextureObject_t textureObject;
int32_t textureWidth, textureHeight;
{
Timer tim("Prepare Image Texture");
const static uint32_t bytesPerPixel{ 3u };
uint8_t* pHostData;
uint32_t bytesPerScanline;
cudaChannelFormatDesc channelFormatDesc;
cudaResourceDesc resourceDesc{};
cudaTextureDesc textureDesc{};
int32_t componentsPerPixel = bytesPerPixel;
pHostData = stbi_load(textureFilename.c_str(), &textureWidth, &textureHeight, &componentsPerPixel, componentsPerPixel);
if (nullptr == pHostData) {
std::cerr << "ERROR: Could not load texture image file '" << textureFilename << std::endl;
return;
}
bytesPerScanline = bytesPerPixel * textureWidth;
channelFormatDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsigned);
checkCudaErrors(cudaMallocArray(&cuArray, &channelFormatDesc, bytesPerScanline, textureHeight));
checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, pHostData, bytesPerScanline * textureHeight * sizeof(uint8_t), cudaMemcpyHostToDevice));
resourceDesc.resType = cudaResourceTypeArray;
resourceDesc.res.array.array = cuArray;
textureDesc.normalizedCoords = true;
textureDesc.filterMode = cudaFilterModePoint;
textureDesc.addressMode[0] = cudaAddressModeWrap;
textureDesc.addressMode[1] = cudaAddressModeWrap;
textureDesc.readMode = cudaReadModeElementType;
checkCudaErrors(cudaCreateTextureObject(&textureObject, &resourceDesc, &textureDesc, nullptr));
STBI_FREE(pHostData);
}
The output images were ppm and converted to pgm for upload purposes