Para o registro, esta é a lição de casa, então ajude tão pouco ou tanto com isso em mente. Estamos usando memória constante para armazenar uma "matriz de máscara" que será usada para realizar uma convolução em uma matriz maior. Quando estou no código do host, estou copiando a máscara para a memória constante usando o cudaMemcpyToSymbol ().

Minha pergunta é uma vez que isso é copiado e eu lanço o código do kernel do dispositivo como o dispositivo sabe onde acessar a matriz de máscara de memória constante. Existe um ponteiro que eu preciso passar no lançamento do kernel. A maior parte do código que o professor nos deu não deve ser alterado (não há nenhum ponteiro para a máscara passada), mas há sempre a possibilidade de ele cometer um erro (embora seja mais provável que eu entenda algo)

A declaração constante da memória deve ser incluída no arquivo kernel.cu separado?

Estou minimizando o código apenas para mostrar as coisas que têm a ver com a memória constante. Como tal, por favor, não aponte se algo não for inicializado ect. Há código para isso, mas isso não é motivo de preocupação no momento.

main.cu:

#include <stdio.h>
#include "kernel.cu"

__constant__ float M_d[FILTER_SIZE * FILTER_SIZE];

int main(int argc, char* argv[])
{

     Matrix M_h, N_h, P_h; // M: filter, N: input image, P: output image

    /* Allocate host memory */
    M_h = allocateMatrix(FILTER_SIZE, FILTER_SIZE);
    N_h = allocateMatrix(imageHeight, imageWidth);
    P_h = allocateMatrix(imageHeight, imageWidth);

    /* Initialize filter and images */
    initMatrix(M_h);
    initMatrix(N_h);


    cudaError_t cudda_ret = cudaMemcpyToSymbol(M_d, M_h.elements, M_h.height * M_h.width * sizeof(float), 0, cudaMemcpyHostToDevice);
    //char* cudda_ret_pointer = cudaGetErrorString(cudda_ret);

    if( cudda_ret != cudaSuccess){
        printf("\n\ncudaMemcpyToSymbol failed\n\n");
        printf("%s, \n\n", cudaGetErrorString(cudda_ret));
    }


    // Launch kernel ----------------------------------------------------------
    printf("Launching kernel..."); fflush(stdout);

    //INSERT CODE HERE
    //block size is 16x16
    //              \\\\\\\\\\\\\**DONE**
    dim_grid = dim3(ceil(N_h.width / (float) BLOCK_SIZE), ceil(N_h.height / (float) BLOCK_SIZE));
    dim_block = dim3(BLOCK_SIZE, BLOCK_SIZE);



    //KERNEL Launch

    convolution<<<dim_grid, dim_block>>>(N_d, P_d);

    return 0;
}

kernel.cu:ESTE É ONDE NÃO SEI COMO ACESSAR A MEMÓRIA CONSTANTE.

//__constant__ float M_c[FILTER_SIZE][FILTER_SIZE];

__global__ void convolution(Matrix N, Matrix P)
{
    /********************************************************************
    Determine input and output indexes of each thread
    Load a tile of the input image to shared memory
    Apply the filter on the input image tile
    Write the compute values to the output image at the correct indexes
    ********************************************************************/

    //INSERT KERNEL CODE HERE

    //__shared__ float N_shared[BLOCK_SIZE][BLOCK_SIZE];


    //int row = (blockIdx.y * blockDim.y) + threadIdx.y;
    //int col = (blockIdx.x * blockDim.x) + threadIdx.x;

}