8
\$\begingroup\$

I wrote this so I can learn CUDA. This is coded to work on my laptop's Nvidia GeForce GT 540M.

Main points I need reviewed:

  1. CUDA programming conventions
  2. Performance, especially kernel speed
  3. C programming conventions
  4. Preprocessor and macro conventions
  5. Making it DRY

kernel.cu

/*************************************************
 *  Part of cuda bitmap to grayscale converter   *
 *  - Bhathiya Perera                            *
 *************************************************/

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#include "utils.cuh"
#include "bitmap.cuh"
#define PIXEL_SIZE 3

cudaError_t turnGrayWithCuda(unsigned char* bitmapData, BitmapInfoHeader* header, unsigned int size);

// Turn given bitmap data to gray scale
__global__ void turnGray(unsigned char* bitmapData, unsigned long size, unsigned int width)
{
#define xIndex (blockIdx.x * blockDim.x + threadIdx.x)
#define yIndex (blockIdx.y * blockDim.y + threadIdx.y)
    unsigned long dataIndex = (xIndex + (yIndex * width)) * PIXEL_SIZE;
#define BLUE bitmapData[dataIndex]
#define GREEN bitmapData[dataIndex+1]
#define RED bitmapData[dataIndex+2]
    // Gray occurs when RED == GREEN == BLUE, so get average
    if(dataIndex < size) {
        unsigned char gray = (RED + GREEN + BLUE) / 3;
        // Convert all pixels to gray
        RED = gray;
        GREEN = gray;
        BLUE = gray;
    }
#undef BLUE
#undef GREEN
#undef RED
#undef xIndex
#undef yIndex
}

void printHelp(char* binary)
{
    printf("GrayScaleCuda\n");
    printf("----------------------------------");
    printf("\t-Bhathiya Perera\n");
    printf("Execute: %s <Bitmap>\n", binary);
}

int main(int argc, char** argv)
{
// Freeing data and calling cudaDeviceReset must be done
// All the time
#undef DO_FAILED_EXIT
#define DO_FAILED_EXIT()\
    free(header);\
    free(data);\
    cudaDeviceReset();\
    return EXIT_FAILURE;

    if (argc != 2) {
        printHelp(argv[0]);
        return EXIT_FAILURE;
    }

#ifdef DEBUG
#define bitmapFilename "C:\\Users\\Bhathiya\\Desktop\\img.bmp"
#else
#define bitmapFilename argv[1]
#endif

    puts("--------------------------------------------------");
    LOG("Welcome to grayscale with cuda.");

    LOG("Turning %s to grayscale...", bitmapFilename);

    BitmapInfoHeader* header = 0;
    header = (BitmapInfoHeader*)malloc(sizeof(BitmapInfoHeader));
    unsigned char* data = loadBitmapFile(bitmapFilename, header);
    if (data==NULL) {
        LOG_ERROR("Failed to load bitmap");
        DO_FAILED_EXIT();
    }

    cudaError_t cudaStatus = turnGrayWithCuda(data, header, header->sizeImage);
    REPORT_CUDA_ERROR(cudaStatus, "Unable to turn grayscale with cuda");

    int success = overwriteBitmapData(bitmapFilename, data);
    if(!success) {
        LOG_ERROR("Failed to overwrite bitmap");
        DO_FAILED_EXIT();
    }

    free(header);
    free(data);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

// Helper function for using CUDA to convert bitmap data to gray
cudaError_t turnGrayWithCuda(unsigned char* bitmapData, BitmapInfoHeader* header, unsigned int size)
{
#undef DO_FAILED_EXIT
#define DO_FAILED_EXIT() cudaFree(devBitmap); return cudaStatus;
    unsigned char* devBitmap = 0;
    cudaError_t cudaStatus;
    size_t dataSize = size * sizeof(unsigned char);
    unsigned long pixelCount = size / PIXEL_SIZE;
    LOG("size=%d, dataSize=%d, pixelCount=%d", size, dataSize, pixelCount);
    LOG("Image Width=%d Height=%d", header->width, header->height);
    cudaStatus = selectBestDevice();
    REPORT_CUDA_ERROR(cudaStatus, "Unable to select a cuda device! "
        "Do you have a CUDA-capable GPU installed?");

    // Allocate GPU buffer for bitmap data
    cudaStatus = cudaMalloc((void**)&devBitmap, dataSize);
    REPORT_CUDA_ERROR(cudaStatus, "Unable allocate device memory");

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(devBitmap, bitmapData, dataSize,
        cudaMemcpyHostToDevice);
    REPORT_CUDA_ERROR(cudaStatus, "Copying memory failed!");

    dim3 threadsPerBlock(32, 32);
    dim3 blocksPerGrid((header->width + threadsPerBlock.x - 1) / threadsPerBlock.x,
        (header->height + threadsPerBlock.y - 1) / threadsPerBlock.y);

    LOG("CUDA kernel launch with %dx%d blocks of %dx%d threads. Total threads=%d",
        blocksPerGrid.x, blocksPerGrid.y, threadsPerBlock.x, threadsPerBlock.y,
        blocksPerGrid.x * blocksPerGrid.y * threadsPerBlock.x * threadsPerBlock.y);

    CUDA_START_TIMER;
    // Launch a kernel on the GPU
    turnGray<<<blocksPerGrid, threadsPerBlock>>>(devBitmap, size, header->width);
    CUDA_STOP_TIMER;

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    REPORT_CUDA_ERROR(cudaStatus, "Kernel launch failed: %s",
        cudaGetErrorString(cudaStatus));

    // Function cudaDeviceSynchronize waits for the kernel to finish, and returns
    // Any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    REPORT_CUDA_ERROR(cudaStatus, "cudaDeviceSynchronize() returned error"
        " code %d after launching kernel!", cudaStatus);

    CUDA_LOG_TIME(size*2.0f/milliseconds/1e6f);

    // Copy bitmap data from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(bitmapData, devBitmap, dataSize,
        cudaMemcpyDeviceToHost);
    REPORT_CUDA_ERROR(cudaStatus, "Copying memory failed!");

    cudaFree(devBitmap);
    return cudaStatus;
}

Additional Content

You can review these too if you want. But I'm only concerned about kernel.cu.

utils.cuh

#ifndef UTILS_CUH
#define UTILS_CUH

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define LOG_ERROR(message, ...) fprintf(stderr, "[ERROR] " message "\n", ##__VA_ARGS__);
#define LOG(message, ...) printf("[INFO] " message "\n", ##__VA_ARGS__);

// Note: `do {} while (0)` is a best practice
// That prevents generating broken code

// Print an array
#define PRINT_ARRAY(data, size)\
    do {\
        LOG("Printing %s...", #data);\
        for(int i = 0; i < size; ++i) {\
            LOG("%s[%02d]=%d", #data, i, data[i]);\
        }\
    } while (0)

#define DO_FAILED_EXIT()

// Make sure that a custom DO_FAILED_EXIT is defined
#define REPORT_CUDA_ERROR(cudaStatus, message, ...)\
    do {\
        if(cudaStatus != cudaSuccess) {\
            fprintf(stderr, "[CUDA-ERROR] " message "\n", ##__VA_ARGS__);\
            DO_FAILED_EXIT();\
        }\
    } while (0)

#ifndef EXIT_SUCCESS
#define EXIT_SUCCESS 0
#endif

#ifndef EXIT_FAILURE
#define EXIT_FAILURE 1
#endif EXIT_FAILURE

#ifndef TRUE
#define TRUE 1
#endif

#ifndef FALSE
#define FALSE 0
#endif

// Note: these two cannot be put into a `do {} while(0)`
// Since `start` and `stop` should be accessible
#define CUDA_START_TIMER\
    cudaEvent_t start, stop;\
    cudaEventCreate(&start);\
    cudaEventCreate(&stop);\
    cudaEventRecord(start);

#define CUDA_STOP_TIMER cudaEventRecord(stop);

#define CUDA_LOG_TIME(bandwidth)\
    do {\
        cudaEventSynchronize(stop);\
        float milliseconds = 0;\
        cudaEventElapsedTime(&milliseconds, start, stop);\
        LOG("Time taken: %f ms", milliseconds);\
        LOG("Effective Bandwidth: %f GB/s", bandwidth);\
    } while (0)


// Select best GPU for processing
cudaError_t selectBestDevice(void);

#endif // UTILS_CUH

utils.cu

#include "utils.cuh"

// Choose which Best Device to run on
// Based on number of multiProcessors available
cudaError_t selectBestDevice(void)
{
    int selectedDevice = 0;
    int numDevices;
    cudaGetDeviceCount(&numDevices);

    if(numDevices > 1) {
        int maxMultiProc = 0;
        int device;
        for(device=0; device < numDevices; ++device) {
            cudaDeviceProp prop;
            cudaGetDeviceProperties(&prop, device);
            if (maxMultiProc < prop.multiProcessorCount) {
                maxMultiProc = prop.multiProcessorCount;
                selectedDevice = device;
            }
        }
    }

    return cudaSetDevice(selectedDevice);
}

bitmap.cuh

#ifndef BITMAP_CUH
#define BITMAP_CUH

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

// Structure packing
#pragma pack(push, 1)

// Define custom types if not defined before
#ifndef WORD
#define WORD short
#endif

#ifndef DWORD
#define DWORD int
#endif

#ifndef LONG
#define LONG unsigned
#endif

// Based on http://stackoverflow.com/questions/14279242/read-bitmap-file-into-structure
// Bitmap structures
typedef struct
{
    WORD type;
    DWORD size;
    WORD reserved1;
    WORD reserved2;
    DWORD offBits;
} BitmapFileHeader;

typedef struct
{
    DWORD size;
    LONG width;
    LONG height;
    WORD planes;
    WORD bitCount;
    DWORD compression;
    DWORD sizeImage;
    LONG xPelsPerMeter;
    LONG yPelsPerMeter;
    DWORD clrUsed;
    DWORD clrImportant;
} BitmapInfoHeader;

#pragma pack(pop)

unsigned char* loadBitmapFile(char* filename, BitmapInfoHeader* bitmapInfoHeader);
int overwriteBitmapData(char* filename, unsigned char* data);

#endif // BITMAP_CUH

bitmap.cu

#include <stdio.h>
#include <stdlib.h>

#include "bitmap.cuh"
#include "utils.cuh"

// Based on http://stackoverflow.com/questions/14279242/read-bitmap-file-into-structure
// Load bitmap file, create a header object and pass it into access information
unsigned char *loadBitmapFile(char* filename, BitmapInfoHeader* bitmapInfoHeader)
{
    FILE* filePtr;
    BitmapFileHeader bitmapFileHeader;
    unsigned char* bitmapImage = 0;
    int imageIdx = 0;

    // Open filename in read binary mode
    filePtr = fopen(filename, "rb");
    if (filePtr == NULL) {
        LOG_ERROR("Failed to open file: '%s'", filename);
        return NULL;
    }

    // Read the bitmap file header
    fread(&bitmapFileHeader, sizeof(BitmapFileHeader), 1, filePtr);

    // Verify that this is a bmp file by checking bitmap id
    if (bitmapFileHeader.type !=0x4D42) {
        LOG_ERROR("Invalid file '%s'", filename);
        fclose(filePtr);
        return NULL;
    }

    // Read the bitmap info header
    fread(bitmapInfoHeader, sizeof(BitmapInfoHeader), 1, filePtr);

    // Move file point to the beginning of bitmap data
    int failed = fseek(filePtr, bitmapFileHeader.offBits, SEEK_SET);

    if(failed != 0) {
        LOG_ERROR("Seeking bitmap failed");
        free(bitmapImage);
        fclose(filePtr);
        return NULL;
    }

    // Read bitmap image data
    bitmapImage = (unsigned char*)malloc(bitmapInfoHeader->sizeImage);

    if (!bitmapImage) {
        LOG_ERROR("Malloc failed");
        free(bitmapImage);
        fclose(filePtr);
        return NULL;
    }

    fread(bitmapImage, bitmapInfoHeader->sizeImage, 1, filePtr);

    // Make sure bitmap image data was read
    if (bitmapImage == NULL) {
        LOG_ERROR("Unable to read bitmap data");
        fclose(filePtr);
        return NULL;
    }

    fclose(filePtr);
    return bitmapImage;
}

int overwriteBitmapData(char* filename, unsigned char* data)
{
    FILE* filePtr;
    BitmapFileHeader bitmapFileHeader;
    BitmapInfoHeader bitmapInfoHeader;

    // Open filename in read write binary mode
    filePtr = fopen(filename, "rb+");
    if (filePtr == NULL) {
        LOG_ERROR("Failed to open file: '%s'", filename);
        return FALSE;
    }

    // Read the bitmap file header
    fread(&bitmapFileHeader, sizeof(BitmapFileHeader), 1, filePtr);

    // Verify that this is a bmp file by checking bitmap id
    if (bitmapFileHeader.type !=0x4D42) {
        LOG_ERROR("Invalid file '%s'", filename);
        fclose(filePtr);
        return FALSE;
    }

    // Read the bitmap info header
    fread(&bitmapInfoHeader, sizeof(BitmapInfoHeader), 1, filePtr);

    // Move file point to the beginning of bitmap data
    int failed = fseek(filePtr, bitmapFileHeader.offBits, SEEK_SET);
    if(failed != 0) {
        LOG_ERROR("Seeking bitmap failed");
        fclose(filePtr);
        return NULL;
    }
    // Overwrite data
    fwrite((void*)data, bitmapInfoHeader.sizeImage, 1, filePtr); 

    fclose(filePtr);
    return TRUE;
}

Output

This is the output I get when I run this program for three different files:

  1. imgSmall.bmp - 56KB
  2. img.bmp - 65MB
  3. imgBig.bmp - 436MB
--------------------------------------------------
[INFO] Welcome to grayscale with cuda.
[INFO] Turning C:\Users\Bhathiya\Desktop\imgSmall.bmp to grayscale...
[INFO] size=57456, dataSize=57456, pixelCount=19152
[INFO] Image Width=167 Height=114
[INFO] CUDA kernel launch with 6x4 blocks of 32x32 threads. Total threads=24576
[INFO] Time taken: 0.027808 ms
[INFO] Effective Bandwidth: 4.132336 GB/s
--------------------------------------------------
[INFO] Welcome to grayscale with cuda.
[INFO] Turning C:\Users\Bhathiya\Desktop\img.bmp to grayscale...
[INFO] size=68623200, dataSize=68623200, pixelCount=22874400
[INFO] Image Width=5400 Height=4236
[INFO] CUDA kernel launch with 169x133 blocks of 32x32 threads. Total threads=23016448
[INFO] Time taken: 21.931168 ms
[INFO] Effective Bandwidth: 6.258052 GB/s
--------------------------------------------------
[INFO] Welcome to grayscale with cuda.
[INFO] Turning C:\Users\Bhathiya\Desktop\imgBig.bmp to grayscale...
[INFO] size=457488000, dataSize=457488000, pixelCount=152496000
[INFO] Image Width=18000 Height=8472
[INFO] CUDA kernel launch with 563x265 blocks of 32x32 threads. Total threads=152775680
[INFO] Time taken: 143.047546 ms
[INFO] Effective Bandwidth: 6.396307 GB/s
\$\endgroup\$
1
6
\$\begingroup\$
  • Just to be really picky: the output should have "CUDA" and not "cuda" since it's an acronym.

  • Also minor: printHelp()'s argument can be a const char* instead.

  • You shouldn't cast the return type of malloc() in C. This is already done for you.

  • I know that macros are common in C, but they still don't have to be overused.

    Here's an example:

    #define xIndex (blockIdx.x * blockDim.x + threadIdx.x)
    #define yIndex (blockIdx.y * blockDim.y + threadIdx.y)
    

    These computations are already very common in CUDA, especially since they're used in place of raw loops. Although they may be a bit long, there's still no harm in typing them all out wherever needed. Plus, if there's any chance for the reader to misunderstand (or forget) exactly what xIndex and yIndex refer to, it could cause problems.

    Here's another example:

    #ifndef EXIT_SUCCESS
    #define EXIT_SUCCESS 0
    #endif
    
    #ifndef EXIT_FAILURE
    #define EXIT_FAILURE 1
    #endif EXIT_FAILURE
    
    #ifndef TRUE
    #define TRUE 1
    #endif
    
    #ifndef FALSE
    #define FALSE 0
    #endif
    

    I don't think there's a chance of them being entirely undefined in a bare CUDA implementation, so you should still be able to use the respective libraries: <stdlib.h> and <stdbool.h>.

    There are other examples of this, but I won't mention them all. Just consider other options, or just omit them, wherever possible.

\$\endgroup\$
6
\$\begingroup\$

I also thought how I could improve my code further.

Readability

First Convert, this

#define xIndex (blockIdx.x * blockDim.x + threadIdx.x)
#define yIndex (blockIdx.y * blockDim.y + threadIdx.y)
    unsigned long dataIndex = (xIndex + (yIndex * width)) * PIXEL_SIZE;
#define BLUE bitmapData[dataIndex]
#define GREEN bitmapData[dataIndex+1]
#define RED bitmapData[dataIndex+2]

to

#define xIndex (blockIdx.x * blockDim.x + threadIdx.x)
#define yIndex (blockIdx.y * blockDim.y + threadIdx.y)
#define BLUE bitmapData[dataIndex]
#define GREEN bitmapData[dataIndex+1]
#define RED bitmapData[dataIndex+2]
    unsigned long dataIndex = (xIndex + (yIndex * width)) * PIXEL_SIZE;

Also, this section needs more comments, and I should avoid magic numbers

dim3 threadsPerBlock(32, 32);
dim3 blocksPerGrid((header->width + threadsPerBlock.x - 1) / threadsPerBlock.x,
    (header->height + threadsPerBlock.y - 1) / threadsPerBlock.y);

Becomes,

// Calculate number of threadsPerBlock and blocksPerGrid
dim3 threadsPerBlock(THREAD_PER_2D_BLOCK, THREAD_PER_2D_BLOCK);
// Need to consider integer devision, and It's lack of precision
// This way total number of threads are newer lower than pixelCount
dim3 blocksPerGrid((header->width + threadsPerBlock.x - 1) / threadsPerBlock.x,
    (header->height + threadsPerBlock.y - 1) / threadsPerBlock.y);

Also this is not formatted correctly

CUDA_LOG_TIME(size*2.0f/milliseconds/1e6f);

Becomes,

// Log Effective Bandwidth and total time
// It is necessary to multiply by 2 because both read and write operations
// Occur
CUDA_LOG_TIME(size * 2.0f / milliseconds / 1e6f);

Performance

Things to remember:

  1. blockDim contains 32 for both x and y since that is what we are using for threadsPerBlock.

  2. Bitwise calculations (shifting <<, >>) are faster than / and *

  3. Instead of multiplying by 32 why not shift left 5 times. (2^5 = 32)

  4. Average of 3 elements can be calculated using averages of 2 elements twice. Also instead of dividing by 2 to calculate average using this method, one can use shift right by 1 time. (2^1 = 2)

Finally

#define PIXEL_SIZE 3
// My GPU Has 1024 Threads per block, thus 32x32 threads
// 32 = 2^5, therefore 32 is 5 bits
#define THREAD_PER_2D_BLOCK 32
#define THREAD_PER_2D_BLOCK_BITS 5

// Turn given bitmap data to gray scale
__global__ void turnGray(unsigned char* bitmapData, unsigned long size, unsigned int width)
{
    // This is done because shifting left by 5 is faster than multiplying by 32
    // Since the 5 least significant bits are zeros after shifting it is possible to use
    // Much faster `bitwise or` to calculate addition. (Since no need to carry the one)
    // The reason it becomes feasible is because threadId is between 0..31
#define xIndex ((blockIdx.x << THREAD_PER_2D_BLOCK_BITS) | threadIdx.x)
#define yIndex ((blockIdx.y << THREAD_PER_2D_BLOCK_BITS) | threadIdx.y)
#define BLUE bitmapData[dataIndex]
#define GREEN bitmapData[dataIndex+1]
#define RED bitmapData[dataIndex+2]
    unsigned long dataIndex = (xIndex + (yIndex * width)) * PIXEL_SIZE;
    // Gray occurs when RED == GREEN == BLUE, so get average
    if(dataIndex < size) {
        // This is done because shifting right is faster than division
        // And average can be calculated in two steps
        unsigned char gray = (((RED + GREEN) >> 1) + BLUE) >> 1;
        // Convert all pixels to gray
        RED = gray;
        GREEN = gray;
        BLUE = gray;
    }
#undef RED
#undef GREEN
#undef BLUE
#undef yIndex
#undef xIndex
}
--------------------------------------------------
[INFO] Welcome to grayscale with CUDA.
[INFO] Turning C:\Users\Bhathiya\Desktop\imgSmall.bmp to grayscale...
[INFO] size=57456, dataSize=57456, pixelCount=19152
[INFO] Image Width=167 Height=114
[INFO] CUDA kernel launch with 6x4 blocks of 32x32 threads. Total threads=24576
[INFO] Time taken: 0.025888 ms
[INFO] Effective Bandwidth: 4.438814 GB/s
--------------------------------------------------
[INFO] Welcome to grayscale with CUDA.
[INFO] Turning C:\Users\Bhathiya\Desktop\img.bmp to grayscale...
[INFO] size=68623200, dataSize=68623200, pixelCount=22874400
[INFO] Image Width=5400 Height=4236
[INFO] CUDA kernel launch with 169x133 blocks of 32x32 threads. Total threads=23016448
[INFO] Time taken: 20.712543 ms
[INFO] Effective Bandwidth: 6.626245 GB/s
--------------------------------------------------
[INFO] Welcome to grayscale with CUDA.
[INFO] Turning C:\Users\Bhathiya\Desktop\imgBig.bmp to grayscale...
[INFO] size=457488000, dataSize=457488000, pixelCount=152496000
[INFO] Image Width=18000 Height=8472
[INFO] CUDA kernel launch with 563x265 blocks of 32x32 threads. Total threads=152775680
[INFO] Time taken: 135.375641 ms
[INFO] Effective Bandwidth: 6.758793 GB/s

Effective Bandwidth Increment: 6.396307 GB/s → 6.758793 GB/s

\$\endgroup\$
2
  • 1
    \$\begingroup\$ The optimizing compiler may already do some of these performance things for you, but yeah, it's still good to be aware of them. \$\endgroup\$ – Jamal Mar 27 '15 at 19:40
  • \$\begingroup\$ @Jamal Yes you are correct. However it shows a performance improvement. I'm using MSVC at the moment. \$\endgroup\$ – bhathiya-perera Mar 27 '15 at 19:49

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

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