#include #include #define PI 3.1415926536 #define PI2 6.283153072 __global__ void _gs_render_cuda( const float *sigmas, const float *coords, const float *colors, float *rendered_img, const int s, // gs num const int h, const int w, const int c, const float dmax ){ int curs = blockIdx.x*blockDim.x + threadIdx.x; if(curs >= s){ return; } float sigma_x = sigmas[curs*3+0]; float sigma_y = sigmas[curs*3+1]; float rho = sigmas[curs*3+2]; float x = coords[curs*2+0]; float y = coords[curs*2+1]; float r = colors[curs*3]; float g = colors[curs*3+1]; float b = colors[curs*3+2]; float negative_half_one_div_one_minus_rho2 = -0.5 / (1-rho*rho); float one_div_sigma_x_2 = 1.0 / sigma_x / sigma_x; float one_div_sigma_y_2 = 1.0 / sigma_y / sigma_y; float two_rho_div_sigma_x_one_div_sigma_y = 2*rho / sigma_x / sigma_y; for(int hi=0; hi dmax || d_y < -dmax){ continue; } for(int wi=0; wi dmax || d_x < -dmax){ continue; } float v = one_div_sigma_x_2*d_x*d_x; v -= two_rho_div_sigma_x_one_div_sigma_y*d_x*d_y; v += one_div_sigma_y_2*d_y*d_y; v *= negative_half_one_div_one_minus_rho2; v = exp(v); atomicAdd(&rendered_img[(hi*w+wi)*c+0], v*r); atomicAdd(&rendered_img[(hi*w+wi)*c+1], v*g); atomicAdd(&rendered_img[(hi*w+wi)*c+2], v*b); } } } void _gs_render( const float *sigmas, const float *coords, const float *colors, float *rendered_img, const int s, const int h, const int w, const int c, const float dmax ) { int threads=64; dim3 grid(int(s/threads)+1); dim3 block(threads); _gs_render_cuda<<>>(sigmas, coords, colors, rendered_img, s, h, w, c, dmax); } __global__ void _gs_render_backward_cuda( const float *sigmas, const float *coords, const float *colors, const float *grads, float *grads_sigmas, float *grads_coords, float *grads_colors, const int s, // gs num const int h, const int w, const int c, const float dmax ){ int curs = blockIdx.x*blockDim.x + threadIdx.x; if(curs >= s){ return ; } // obtain parameters of gs float sigma_x = sigmas[curs*3+0]; float sigma_y = sigmas[curs*3+1]; float rho = sigmas[curs*3+2]; float x = coords[curs*2+0]; float y = coords[curs*2+1]; // float w1 = -0.5 / (1-rho*rho) ; float w2 = 1.0 / (sigma_x*sigma_x); float w3 = 1.0 / (sigma_x*sigma_y); float w4 = 1.0 / (sigma_y*sigma_y); float od_sx = 1.0 / sigma_x; float od_sy = 1.0 / sigma_y; // init for(int hi = 0; hi < h; hi++){ for( int wi=0; wi < w; wi++){ float curw_f = 2.0*wi/(w-1) - 1.0; float curh_f = 2.0*hi/(h-1) - 1.0; // compute the 2d gs value float d_x = curw_f - x; // distance along x axis float d_y = curh_f - y; if(d_x > dmax || d_x < -dmax || d_y > dmax || d_y < -dmax){ continue; } float d = w2*d_x*d_x - 2*rho*w3*d_x*d_y + w4*d_y*d_y; float v = w1*d; v = exp(v); // printf("si:%d, sigma_x: %f, sigma_y:%f, rho:%f, x:%f, y:%f, v:%f\n", si, sigma_x, sigma_y, rho, x,y,v); // compute grad of coords float v_2_w1 = v*2*w1; float g_vst_to_gsx = v_2_w1*(-w2*d_x+rho*w3*d_y); // grad of v^{st} to G^s_x float g_vst_to_gsy = v_2_w1*(-w4*d_y+rho*w3*d_x); // grad of v^{st} to G^s_y // compute grad of sigmas float g_vst_to_gsigx = v_2_w1*od_sx* (w3*rho*d_x*d_y - w2*d_x*d_x); float g_vst_to_gsigy = v_2_w1*od_sy* (w3*rho*d_x*d_y - w4*d_y*d_y); float g_vst_to_rho = -v_2_w1*(2*w1*rho*d+w3*d_x*d_y); for(int ci=0; ci>>(sigmas, coords, colors, grads, grads_sigmas, grads_coords, grads_colors, s, h, w, c, dmax); }