Q:

Allocation of variables in CUDA C++ - unable to allocate device variable

I'm trying to process .bmp file with filter on GPU. For this I need to allocate and fill variable d_bufferRGB on device. But when I do so, my compiler just tells me it is defined on host and end with error "identifier d_bufferRGB is undefined". Functions load_BMP and save_BMP are fully functional. I need to know how to copy bufferRGB into d_bufferRGB, cause obviously I am doing some mistake.
Can you please help me ?
The source code is here >>>
  #include <stdio.h>
#include <stdlib.h>
#include <Windows.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include "device_launch_parameters.h"
#include <iostream>
//using namespace std;

int width, heigth;
long size, *d_size;
RGBTRIPLE *bufferRGB, *new_bufferRGB;
RGBTRIPLE *d_bufferRGB, *d_new_bufferRGB;

void load_bmp(RGBTRIPLE **bufferRGB, int *width, int *heigth, const char *file_name)
{
    BITMAPFILEHEADER bmp_file_header;
    BITMAPINFOHEADER bmp_info_header;
    FILE *file;

    file = fopen(file_name, "rb");

    fread(&bmp_file_header, sizeof(BITMAPFILEHEADER), 1, file);

    fread(&bmp_info_header, sizeof(BITMAPINFOHEADER), 1, file);

    *width = bmp_info_header.biWidth;
    *heigth = bmp_info_header.biHeight;
    size = (bmp_file_header.bfSize - bmp_file_header.bfOffBits);

    int x, y;
    *bufferRGB = (RGBTRIPLE *)malloc(*width* *heigth * 4);

    fseek(file, bmp_file_header.bfOffBits - sizeof(bmp_file_header) - sizeof(bmp_info_header), SEEK_CUR);

    for (y = 0; y < *heigth; y++)
    {
        for (x = 0; x < *width; x++)
        {
            (*bufferRGB)[(y * *width + x)].rgbtBlue = fgetc(file);
            (*bufferRGB)[(y * *width + x)].rgbtGreen = fgetc(file);
            (*bufferRGB)[(y * *width + x)].rgbtRed = fgetc(file);
        }
        for (x = 0; x < (4 - (3 * *width) % 4) % 4; x++)
            fgetc(file);
    }
    fclose(file);
}

void save_bmp(RGBTRIPLE *bufferRGB, const char *new_name, const char *old_name)
{
    BITMAPFILEHEADER bmp_file_header;
    BITMAPINFOHEADER bmp_info_header;
    FILE *file;

    file = fopen(old_name, "rb");

    fread(&bmp_file_header, sizeof(BITMAPFILEHEADER), 1, file);

    fread(&bmp_info_header, sizeof(BITMAPINFOHEADER), 1, file);
    fclose(file);

    file = fopen(new_name, "wb");

    fwrite(&bmp_file_header, sizeof(BITMAPFILEHEADER), 1, file);
    fwrite(&bmp_info_header, sizeof(BITMAPINFOHEADER), 1, file);
    fseek(file, bmp_file_header.bfOffBits - sizeof(bmp_file_header) - sizeof(bmp_info_header), SEEK_CUR);

    int alligment_x = (4 - (3 * width) % 4) % 4;
    unsigned char *to_save = (unsigned char *)malloc((width * 3 + alligment_x)*heigth);
    unsigned int index = 0;
    int x, y;

    for (y = 0; y < heigth; y++)
    {
        for (x = 0; x < width; x++)
        {
            to_save[index++] = bufferRGB[(y * width + x)].rgbtBlue;
            to_save[index++] = bufferRGB[(y * width + x)].rgbtGreen;
            to_save[index++] = bufferRGB[(y * width + x)].rgbtRed;
        }
        for (x = 0; x < alligment_x; x++)
            to_save[index++] = 0;
    }
    fwrite(to_save, (width * 3 + alligment_x)*heigth, 1, file);

    fclose(file);
    free(to_save);
}

__global__ void CUDA_filter_grayscale()
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    BYTE grayscale;

    if (idx < *d_size)
    {
        grayscale = ((d_bufferRGB[idx].rgbtRed + d_bufferRGB[idx].rgbtGreen + d_bufferRGB[idx].rgbtBlue) / 3);
        d_new_bufferRGB[idx].rgbtRed = grayscale;
        d_new_bufferRGB[idx].rgbtGreen = grayscale;
        d_new_bufferRGB[idx].rgbtBlue = grayscale;
    }
}

int main()
{

    //GPU cast ----------------------------------------------------------------------------------------------

    cudaMalloc((void**)&d_new_bufferRGB, width*heigth * 4);
    cudaMalloc((void**)&d_bufferRGB, width*heigth * 4);
    cudaMalloc((void**)&d_size, sizeof(size));

    load_bmp(&bufferRGB, &width, &heigth, "test.bmp");

    cudaMemcpy(&d_size, &size, sizeof(size), cudaMemcpyHostToDevice);
    cudaMemcpy(d_bufferRGB, bufferRGB, sizeof(bufferRGB), cudaMemcpyHostToDevice);

    CUDA_filter_grayscale << <256, 512 >> >(); //size dont bother me for now

    cudaMemcpy(d_new_bufferRGB, new_bufferRGB, sizeof(new_bufferRGB), cudaMemcpyDeviceToHost);

    save_bmp(new_bufferRGB, "filter_grayscale_GPU", "test.bmp");

}

A:

Kernels cannot access host-code variables (such as the variable d_bufferRGB). If you need access to a host-code variable in a kernel, you need to pass it to the kernel as an argument:
__global__ void CUDA_filter_grayscale(const RGBTRIPLE *d_bufferRGB, RGBTRIPLE *d_new_bufferRGB)
{
    // Code same as before
}

// Usage in main():
CUDA_filter_grayscale << <256, 512 >> >(d_bufferRGB, d_new_bufferRGB);

