Convert a 24bit bitmap to grayscale

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:

• CUDA programming conventions

• Performance, especially kernel speed

• C programming conventions

• Preprocessor and macro conventions

• Making it DRY

kernel.cu

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

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

#include
#include

#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 \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));
unsig

-
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: ` and `.

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

StackExchange Code Review Q#85091, answer score: 6