Para el registro, esto es tarea, así que ayude tan poco o mucho con eso en mente. Estamos utilizando memoria constante para almacenar una "matriz de máscara" que se utilizará para realizar una convolución en una matriz más grande. Cuando estoy en el código de host, estoy copiando la máscara a la memoria constante usando el cudaMemcpyToSymbol ().

Mi pregunta es una vez que se haya copiado y lancé el código del kernel de mi dispositivo. ¿Cómo sabe el dispositivo dónde acceder a la matriz de máscara de memoria constante? ¿Hay un puntero que necesito pasar en el lanzamiento del kernel? La mayoría del código que el profesor nos dio no debe cambiarse (no hay un puntero a la máscara que se pasa) pero siempre existe la posibilidad de que cometió un error (aunque es muy probable que mi comprensión de algo)

¿Se supone que la declaración constante de memoria se debe incluir en el archivo kernel.cu separado?

Estoy minimizando el código para mostrar solo las cosas que tienen que ver con la memoria constante. Como tal, por favor no señale si algo no está inicializado, etc. Hay un código para eso, pero eso no es motivo de preocupación en este 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 ES DONDE NO SABE CÓMO ACCEDER A LA MEMORIA 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;

}