Temario

Sesión	Tarea
1. Introducción a la programación en CUDA
2. Ejemplos usando CUDA	Tarea_1_VectoresMatrices.pdf
3. Procesamiento de Imágenes
4. Procesamiento de Imágenes usando CUDA	Tarea_2_CambiarColorImagen.pdf Tarea_3_CombinacionImagenes.pdf Tarea_4_DeteccionBordes.pdf

Calificaciones:
             

Num. Ctrl.	T1	Comentarios	T2	Comentarios	T3	Comentarios	T4	Comentarios	Prom. Tareas	Examen	Proyecto	Calif. Final (0% - 100%) Prog. en CUDA
15216336	10		8	T2_1	10	T3_1	9.5	T4_1, T4_2	9.375	6.0	10	92.5
9003089	9	T1_14, T1_15	10		10	T3_1	10		9.75	7.5	10	98.5
15216330	8.5	T1_4, T1_12	10		10	T3_1	6	T4_3,T4_4,T4_5	8.625	7.0	7	83.5
1234567	9	T1_13	10		10	T3_1	10		9.75	8.0	10	100

    
                                                                          
Comentarios:

• T1_1: Al usar C/C++, cuando creamos memoria a un vector de tamaño N, los indices que recorren al vector van de [0,N-1].
• T1_2: Al mandar a llamar a tu función kernel de esta manera:

          Sum_Vec<<<1,N-1>>>(v_d,s_d);
          Se ejecutan N-1 hilos, entonces dentro de la función kernel, los indices de threadIdx.x van de [0,N-2]
          El detalle en tu codigo es que estas leyendo la información de tu vector V en la localidad 0, pero ahí no pusiste información, entonces este valor puede ser basura o en el mejor caso 0.

• T1_3: El resultado del ejercicio 1b es incorrecto.
• T1_4: En tu ejercicio 2 hay una copia de memoria no necesaria: cudaMemcpy(C_d,C_h,size,cudaMemcpyHostToDevice);
  
• T1_5: Tu ejercicio 2 no compila y las funciones kernel estan mal programadas.
• T1_6: No entregó el reporte.
• T1_7: En el ejercicio 1b, faltó dividir entre dos.
• T1_8: El parámetro alpha solo se pasa como parámetro. No es necesario crear memoria para este parámetro con cudaMalloc, ya que solo es un escalar.
• T1_9: Está mal convertido el valor de entera a float: (float *)x, debe ser: (float )x
• T1_10: Para datos float en el printf hay que poner %f
• T1_11: Tus matrices A_h, B_h y C_h deben ser creadas de forma estática float A_h[FILAS][COLUMNAS], ya que en tu asignación usas: A_h[y][x] =
  
• T1_12: En el ejercicio 2b, la formula está mal implementada. Además, el \alpha debe ser un valor real entre [0,1].
• T1_13: Copias de memoria no necesarias en tus ejercicios:

            cudaMemcpy(B_dev,B,(N-1)*sizeof(float),cudaMemcpyHostToDevice);
            cudaMemcpy(C1_dev,C1,N*N*sizeof(float),cudaMemcpyHostToDevice);
            cudaMemcpy(C2_dev,C2,N*N*sizeof(float),cudaMemcpyHostToDevice);
            cudaMemcpy(C_dev,C,(N-2)*sizeof(float),cudaMemcpyHostToDevice);
            cudaMemcpy(A,A_dev,N*sizeof(float),cudaMemcpyDeviceToHost);
            cudaMemcpy(A,A_dev,N*N*sizeof(float),cudaMemcpyDeviceToHost);
            cudaMemcpy(B,B_dev,N*N*sizeof(float),cudaMemcpyDeviceToHost);

•  T1_14: El ejercicio 2 solo funciona para matrices cuadradas. Ya que tu bloque de hilos esta definido como:
  

            dim3 block(N,M,1) ;
            int x = blockIdx.x*blockDim.x+threadIdx.x;
            int y = blockIdx.y*blockDim.y+threadIdx.y;
            Entonces "x" va sobre las filas desde 0 hasta N-1, y "y" va sobre las columnas desde 0 hasta M-1. Por lo tanto, el indice global int idxA=x*N+y; está mal calculado.

• T1_15: Copias de memoria no necesarias en tus ejercicios:
  

            cudaMemcpy(S1_dev,S1,size_S1,cudaMemcpyHostToDevice);
            cudaMemcpy(S2_dev,S2,size_S2,cudaMemcpyHostToDevice);
            cudaMemcpy(C1_dev,C1,size_C,cudaMemcpyHostToDevice);
            cudaMemcpy(C2_dev,C2,size_C,cudaMemcpyHostToDevice);

• T2_1: Hay varios detalles a considerar en tu programa serial, y no compila:
• La variable frame_Original, la declaraste con O mayúscula y la estas usando con o minúscula.
• El cv::vect3b hay que cambiarlo por cv::Vec3b.
• En la parte de aplicar el umbral, la variable Dst no ha sido inicializada, por lo tanto tu imagen resultante no es la esperada.

• T3_1: Mejora el resultado si aplicas un umbral (valorPixel > 127) a la imagen alpha (greenscreenMask).
  

• T4_1: La imagen en donde guardas el resultado de la MG, tiene el ancho y alto invertidos, por eso al copiar la memoria del Device al Host la imagen no se ve bien. Debes crear la memoria de la siguiente manera: temp.create(imageHeight, imageWidth, CV_32FC(1));
• T4_2: Los filtros es mejor declararlos del tipo: __constant__ float fuera de las funciones Kernel.
• T4_3: Al definir tus bloques y malla de hilos de forma unidimensional:

           dim3 threads(512, 1, 1);
           int N = imageW * imageH;
           dim3 grid(N / threads.x + (N%threads.x == 0 ? 0 : 1), 1, 1);
           No debes usar los índices de los hilos de forma bidimensional dentro de tus funciones Kernel (filtroX y filtroY):
           const int ix= blockDim.x * blockIdx.x + threadIdx.x;
           const int iy= blockDim.y * blockIdx.y + threadIdx.y;

• T4_4: Nota que al aplicar los filtros K1 y K2 de derivadas, es posible tener valores negativos, entonces no es posible usar un tipo de dato unsigned char para guardar el resultado.
• T4_5: El resultado es incorrecto.

Bibliografía:

• Sanders, J., & Kandrot, E. (2010). CUDA by Example: An Introduction to General-Purpose GPU Programming, Portable Documents. Addison-Wesley Professional.
• Kirk, D. B., & Wen-Mei, W. H. (2009). Programming massively parallel processors: a hands-on approach. Morgan Kaufmann.
• Cook, S. (2012). CUDA programming: a developer's guide to parallel computing with GPUs. Newnes.
• Wilt, N. (2013). The cuda handbook: A comprehensive guide to gpu programming. Pearson Education.
• GPU Technology Conference, http://www.gputechconf.com/
• Webinars, https://developer.nvidia.com/gpu-computing-webinars