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 {
    __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;
        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;
    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;
        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));

The output images were ppm and converted to pgm for upload purposes From CPU code

From CUDA code

  • 1
    $\begingroup$ This is probably not anything to do with the texture, and almost certainly nothing to do with double vs float. Look closely at the edge of the sphere: in the first image it is sharp, and in the second blurry, blending into the background color. This suggests to me that it's the ray distribution on the image plane that's the issue. Is the second image distributing rays over a larger area per pixel, using a wider AA filter, etc? $\endgroup$ Commented Nov 22, 2020 at 17:34
  • 1
    $\begingroup$ @NathanReed Thanks for the response. Indeed it turned out to be the camera aperture: on the GPU it was 0.1f whereas on the CPU 0.0f. If you like to provide your input as an "answer", I will gladly mark it so. $\endgroup$
    – Vectorizer
    Commented Dec 2, 2020 at 17:31
  • $\begingroup$ From Ray Tracing in One Weekend:For our virtual camera, we can have a perfect sensor and never need more light, so we only have an aperture when we want defocus blur. $\endgroup$
    – Vectorizer
    Commented Dec 2, 2020 at 18:45

1 Answer 1


If you look closely at the edge of the sphere, in the first image it is sharp, and in the second image it's blurry, blending into the background color. This suggests to me that it's not the texture, but the ray distribution on the image plane that's the issue. The second image might be distributing rays over a larger area per pixel, for instance by using a wider AA filter, or a wider lens aperture (defocus blur).


Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.