Demo: Multiplicação de Matriz por Escalar (Em CUDA)

Multiplicação de F * E = R onde F é uma matriz, E um escalar e R uma matriz.

Kernel + Códigos em C

 1 #include <cuda.h>
 2 #include <string.h>
 3 #include <math.h>
 4 #include "simple_arrays.h"
 5 
 6 #undef _POSIX_C_SOURCE
 7 #undef _XOPEN_SOURCE
 8 
 9 #define wB 32.0
10 #define hB 16.0
11 #define sB wB*hB
12 
13 
14 __global__ void kn_Escalar(float *dF,float *dR,int w , int h, int E){
15 
16     int iX = blockIdx.x*blockDim.x + threadIdx.x;
17     int iY = (blockIdx.y*blockDim.y) + threadIdx.y;
18 
19 
20     if( iX >= w || iY >= h ){ return;}
21 
22     dR[iY*w+iX] = dF[iY*w+iX] * E;
23 
24 }
25 
26 void cd_Escalar( float *F,float *R, int width , int height,int E){
27 
28     float *dF,*dR;
29 
30     int sF = width*height*sizeof(float);
31 
32     cudaMalloc((void **)&dF, sF);
33     cudaMalloc((void **)&dR, sF);
34 
35     cudaMemcpy(dF, F, sF, cudaMemcpyHostToDevice);
36     cudaMemset(dR, 0, sF);
37 
38     dim3 Block ( wB,hB ); // #Define Acima
39     dim3 Grid  ( ceil( width/wB ),ceil( height/hB ) );
40 
41     kn_Escalar<<<Grid,Block>>>(dF,dR,width,height,E);
42     cudaThreadSynchronize();
43 
44     cudaMemcpy(R, dR, sF, cudaMemcpyDeviceToHost);
45 
46     cudaFree(dF);
47     cudaFree(dR);
48 }
49 
50 void Escalar(Image32F *F,int E,Image32F *R)
51 {
52 
53     R->set_dims(F->nd,F->dims);
54     R->raster =(char *)new float[F->size];
55 
56     cd_Escalar( (float*)F->raster,
57                 (float*)R->raster,
58                  F->dims[0],F->dims[1],
59                  E );
60 
61 }
ERROR ping: sandbox "xsb_pycuda" is busy. Try again.

Cabeçalho da Função em C

1 void Escalar(Image32F *in,int e,Image32F *out);
ERROR ping: sandbox "xsb_pycuda" is busy. Try again.

Chamada Python da Função

** Estes valores não podem ser muito grandes para não estourar a imagem de alocação do bloco.
ERROR ping: sandbox "xsb_pycuda" is busy. Try again.
 1 from ia636 import *
 2 from numpy import *
 3 import primo as PR
 4 
 5 
 6 F = arange( H*W ,dtype=float32 )
 7 F = F.reshape(H,W)
 8 
 9 R = PR.Escalar(F,E)
10 
11 print "F:"
12 print F
13 print ""
14 print "E:", E
15 print ""
16 print "R:"
17 print R
ERROR ping: sandbox "xsb_pycuda" is busy. Try again.

Vizualizando a Alocação do Grid

 1 import Image, ImageDraw, ImageFont
 2 
 3 mmfreedom(2)
 4 
 5 def frameit(image, dx, dy):
 6     h, w = image.shape[1:]
 7     res = 255*ones((3, h+2*dy, w+2*dx), image.dtype)
 8     res[:,dy:-dy,dx:-dx] = image
 9     return res
10 
11 def chessboard_image(image, m1=None, m2=None, m3=None, pix_size=25, fill_cell='#dee7ec',
12                                                                     fill_cell2='#ee858a',
13                                                                     font_size = 12):
14     import math
15     h, w = image.shape
16 
17     if m1 is None:
18         m1 = zeros(image.shape, bool)
19     if m2 is None:
20         m2 = zeros(image.shape, bool)
21     if m3 is None:
22         m3 = zeros(image.shape, bool)
23 
24     chessb = 255*ones((3, pix_size*h+1, pix_size*w+1), uint8)
25     chessb[:,::pix_size,:] = 0
26     chessb[:,:,::pix_size] = 0
27 
28     pil = array2pil(chessb)
29     font = ImageFont.truetype('/usr/share/fonts/truetype/msttcorefonts/cour.ttf', font_size)
30 
31     draw = ImageDraw.Draw(pil)
32 
33     for i in range(w):
34         for j in range(h):
35             v = image[j,i]
36             n = 0
37             if v > 0: n = int(math.log10(v))
38             y0 = pix_size*i
39             x0 = pix_size*j
40             y = y0 + [9, 5, 2][n]
41             x = x0 + 7
42             if m2[j,i]:
43                 draw.rectangle((y0+1, x0+1, y0+pix_size-1, x0+pix_size-1), fill=fill_cell)
44             if m3[j,i]:
45                 draw.rectangle((y0+3, x0+3, y0+pix_size-3, x0+pix_size-3), fill=fill_cell2)
46             if m1[j,i]:
47                 draw.text((y,x), '%d' % v, font=font, fill='#000000')
48 
49     chessb = pil2array(pil)
50     return frameit(chessb, 5, 5)
ERROR ping: sandbox "xsb_pycuda" is busy. Try again.

Matrizes F e R, divisão da matriz em um blocos de 32x16 como definido nas linas 38-42.

A linha 20 do kernel impede as threads fora da imagem de executarem.

ERROR ping: sandbox "xsb_pycuda" is busy. Try again.