From: Gilles Perrot Date: Wed, 22 Jun 2011 06:06:05 +0000 (+0200) Subject: version operationnelle avec chemins parametrables X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/lniv_gpu.git/commitdiff_plain/7cc36a5689750ee7d77b959362edc65948a79427 version operationnelle avec chemins parametrables --- diff --git a/.gitignore b/.gitignore index e69de29..c39d6c5 100644 --- a/.gitignore +++ b/.gitignore @@ -0,0 +1,2 @@ +# ne pas surveiller les fichiers objet +*.o \ No newline at end of file diff --git a/Makefile b/Makefile index 2801dda..20361ca 100644 --- a/Makefile +++ b/Makefile @@ -34,9 +34,9 @@ ################################################################################ # Add source files here -EXECUTABLE := levelines +EXECUTABLE := lniv # CUDA source files (compiled with cudacc) -CUFILES := main_gmem.cu +CUFILES := main.cu # CUDA dependency files CU_DEPS := levelines_common.h # C/C++ source files (compiled with gcc / c++) diff --git a/image_out16.pgm b/image_out16.pgm new file mode 100644 index 0000000..2cbfd03 Binary files /dev/null and b/image_out16.pgm differ diff --git a/levelines_common.h b/levelines_common.h index ddb0ea2..b96bc6b 100644 --- a/levelines_common.h +++ b/levelines_common.h @@ -14,6 +14,159 @@ // Reference CPU functions //////////////////////////////////////////////////////////////////////////////// //extern "C" void fonc(...); +/** + * + * \brief determine la valeur de NG de la ligne de niveau demarrant de i,j + * \author NB - PhyTI + * + * \param[in] image image d'entree + * \param[in] i coord i du point de depart + * \param[in] j coord j du point de depart + * + * \return la valeur de la ligne + * + * la ligne de niveau ne va que d'un cote + * on cherche le segment de r pixels ou les valeurs + * de l'image ont une variance minimale + * il y a 24 segments possibles (4 par quadrants) + * la direction 0 va vers le haut, les directions tournent + * dans le sens horaire, dans l'ordre croissant + * + * pas de test de bord : pas d'utilisation a moins de r pixels du bords + * + * + * EN TEST + */ +int lniv4_value(unsigned int *image, int2 *path, int i, int j, int idim, int jdim, int *dout, unsigned int r ) +{ + int value_c ; + int value2_c ; + int d, v, p, d_min, eq_min, eq, sum, sum2 ; + int sum_eq_min ; + int it, jt ; + + + /* mem */ + value_c = image[i*jdim + j] ; + value2_c = value_c*value_c ; + + // direction d=0 + sum = value_c ; + sum2 = value2_c ; + it = i ; + jt = j ; + for (p=0; p à faire disparaître + * longueur = 4 pixels + * une ligne = un chemin + ************************************************************************/ __constant__ int pathDi[PSIZE_I][PSIZE_J-1] = { // Q1 - { -1, -1, -1}, // - { -1, -1, -1}, // - { -1, -1, -1}, // - { -1, -1, -1}, // - { -1, 0, 1}, // - { 0, -1, 0}, + { -1, -1, -1}, // 90 + { -1, -1, -1}, // 75 + { -1, -1, -1}, // 60 + { -1, -1, -1}, // 45 + { -1, 0, -1}, // 30 + { 0, -1, 0}, // 15 // Q4 - { 0, 0, 0}, // - { 0, 1, 1}, // - { 1, 0, 1}, // - { 1, 1, 1}, // - { 1, 1, 1}, // - { 1, 1, 1}, + { 0, 0, 0}, // 0 + { 0, 1, 0}, // 345 + { 1, 0, 1}, // 330 + { 1, 1, 1}, // 315 + { 1, 1, 1}, // 300 + { 1, 1, 1}, // 285 // Q3 - { 1, 1, 1}, // - { 1, 1, 1}, // - { 1, 1, 1}, // - { 1, 1, 1}, // - { 1, 0, -1}, // - { 0, 1, 0}, + { 1, 1, 1}, // 270 + { 1, 1, 1}, // 255 + { 1, 1, 1}, // 240 + { 1, 1, 1}, // 225 + { 1, 0, 1}, // 210 + { 0, 1, 0}, // 195 // Q2 - { 0, 0, 0}, // - { 0, -1, 0}, // - { -1, 0, -1}, // - { -1, -1, -1}, // - { -1, -1, -1}, // - { -1, -1, -1} + { 0, 0, 0}, // 180 + { 0, -1, 0}, // 165 + { -1, 0, -1}, // 150 + { -1, -1, -1}, // 135 + { -1, -1, -1}, // 120 + { -1, -1, -1} // 105 } ; // __constant__ int pathDj[PSIZE_I][PSIZE_J-1] = { // Q1 - { 0, 0, 0}, // - { 0, 1, 0}, - { 1, 0, 1}, - { 1, 1, 1}, - { 1, 1, 1}, - { 1, 1, 1}, + { 0, 0, 0}, // 90 + { 0, 1, 0}, // 75 + { 1, 0, 1}, // 60 + { 1, 1, 1}, // 45 + { 1, 1, 1}, // 30 + { 1, 1, 1}, // 15 // Q4 - { 1, 1, 1}, // - { 1, 1, 1}, - { 1, 1, 1}, - { 1, 1, 1}, - { 1, 0, -1}, - { 0, 1, 0}, - // Q3 - { 0, 0, 0}, // - { 0, -1, 0}, - { -1, 0, -1}, - { -1, -1, -1}, - { -1, -1, -1}, - { -1, -1, -1}, + { 1, 1, 1}, // 0 + { 1, 1, 1}, // 345 + { 1, 1, 1}, // 330 + { 1, 1, 1}, // 315 + { 1, 0, 1}, // 300 + { 0, 1, 0}, // 285 + // Q3 + { 0, 0, 0}, // 270 + { 0, -1, 0}, // 255 + { -1, 0, -1}, // 240 + { -1, -1, -1}, // 225 + { -1, -1, -1}, // 210 + { -1, -1, -1}, // 195 // Q2 - { -1, -1, -1}, // - { -1, -1, -1}, - { -1, -1, -1}, - { -1, -1, -1}, - { -1, 0, 1}, - { 0, -1, 0} + { -1, -1, -1}, // 180 + { -1, -1, -1}, // 165 + { -1, -1, -1}, // 150 + { -1, -1, -1}, // 135 + { -1, 0, -1}, // 120 + { 0, -1, 0} // 105 } ; + + +// valeurs des tangentes des angles de base pour la génération initiale des chemins +// pour la version à chemins de longueur paramétrable +__constant__ float tangente[] = {0.000, 0.268, 0.577, 1.000} ; -// declare texture reference for 2D int texture +// declarations des textures texture tex_img_in ; texture tex_img_estim ; texture tex_img_lniv ; +texture tex_paths ; + + + +/** + * + * \brief calcule les chemins + * \author NB - PhyTI, modifié by zulu pour adaptater aux chemins paramétrables + * + * \param[in] r longueur des chemins + * + * \param[out] d_paths matrice des déplacements relatifs (chemins) + * + * Cette fonction utilise le tableau constant des tangentes des angles + * considérés pour le calcul de chemins (float tangente[]). + * + */ +__global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){ + + unsigned int idpath = 0 ; + int ic, jc, iprec, jprec ; + float offset = 0.5 ; + unsigned int basepath = 0 ; + + // Q1 inf + for (int a=0 ; a< 4 ; a++){ // les 4 angles 0,15,30 et 45 + for (int p=0 ; p< r ; p++){ // les r points + ic = r-1 - floor(tangente[a]*p + offset) ; + if ( p > 0 ){ + d_paths[idpath*(r-1)+p-1].x = ic - iprec ; + d_paths[idpath*(r-1)+p-1].y = 1 ; + } + iprec = ic ; + } + idpath++ ; + } + // Q1 sup + for (int a=2 ; a>0 ; a--){ // les 2 angles 60 et 75 + for (int p=0 ; p< r ; p++){ // les r points + jc = floor(tangente[a]*p + offset) ; + if ( p > 0 ){ + d_paths[idpath*(r-1)+p-1].x = -1 ; + d_paths[idpath*(r-1)+p-1].y = jc - jprec ; + } + jprec = jc ; + } + idpath++ ; + } + + // Q2 + basepath += 6 ; + for (int a=0 ; a< 6 ; a++){ // les 6 angles 90,105,120,135,150,165 + for (int p=0 ; plpath)&&(ilpath)&&(jlpath)&&(ilpath)&&(j=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath)){ + z = tex2D(tex_img_estim, j, i) ; for( idpath=0; idpath < PSIZE_I ; idpath++) { ic = i ; jc = j ; - zc = tex2D(tex_img_estim, j, i) ; - mse.x = zc ; - mse.y = zc*zc ; + mse.x = z ; + mse.y = z*z ; for( idpix=0; idpix < lpath-1 ; idpix++ ) { - ic += pathDi[idpath][idpix] ; - jc += pathDj[idpath][idpix] ; + ic += tex2D(tex_paths, idpix, idpath).x ; + jc += tex2D(tex_paths, idpix, idpath).y ; zc = tex2D(tex_img_estim, jc, ic) ; mse.x += zc ; mse.y += zc*zc ; } // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath ) - // a ameliorer pour vitesse + // TODO cherchera ameliorer pour vitesse mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ; - if (idpath == 0) { + if ( (idpath == 0) || (mse_cur < mse_min) ) { mse_min = mse_cur ; - val = mse.x ; - } else { - if ( mse_cur < mse_min ) { - mse_min = mse_cur ; - val = mse.x ; - } + val = mse.x ; } } img_out[ i*L + j ] = val / lpath ; @@ -241,60 +412,101 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L, } +/** + * + * \brief determine les lniv en chaque point de l'image + * \author zulu - AND + * + * \param[in] L Largeur de l'image + * \param[in] H Hauteur de l'image + * \param[in] r longueur des segments + * + * \param[out] img_out image des lniv + * + * Execution sur des blocs de threads 2D et une grille 2D + * selon les dimensions de l'image. + * L'image d'entrée doit être au préalable en mémoire texture pointée par "tex_img_estim". + * Les matrices des chemins sont, elles, pointées par "tex_paths" + * Cette version ne fournit pas les indices des chemins pour les tracé éventuel des lniv. + * Cette version tente d'utiliser la shared memory pour compenser la baisse de perf due aux chemins + * paramétrables non constants. + */ - -__global__ void kernel_levelines_only_texture(unsigned int * img_out, unsigned int L, unsigned int H) +__global__ void kernel_levelines_texture_smem(unsigned int * img_out, unsigned int L, unsigned int H, unsigned int r) { + // coordonnées du point dans le bloc + unsigned int iib = threadIdx.x ; + unsigned int jib = threadIdx.y ; // coordonnes du point dans l'image - unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; - unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;; - //unsigned int spos = threadIdx.x * blockDim.y + threadIdx.y ; - + unsigned int i = blockIdx.x*blockDim.x + iib ; + unsigned int j = blockIdx.y*blockDim.y + jib ; + // nb de points par chemin - int lpath = PSIZE_J ; - unsigned int ic, jc ; + int lpath = r ; + int ic, jc ; int idpath, idpix ; - unsigned int mse_min, mse_cur ; - - //extern __shared__ uint2 mse[] ; - uint2 mse ; + unsigned int val, mse_cur, mse_min, z, zc ; + uint2 mse_data ; + + //__shared__ unsigned int val_img[16*16] ; + + //val_img[jib*16+iib] = tex2D(tex_img_estim, j, i) ; - if((i>lpath)&&(ilpath)&&(j=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath)){ + z = tex2D(tex_img_estim, j, i) ; for( idpath=0; idpath < PSIZE_I ; idpath++) { ic = i ; jc = j ; - mse.x = tex2D(tex_img_in, i, j) ; - mse.y = tex2D(tex_img_in, i, j)*tex2D(tex_img_in, i, j) ; + mse_data.x = z ; + mse_data.y = z*z ; + mse_min = mse_data.y - mse_data.x/lpath*mse_data.y ; for( idpix=0; idpix < lpath-1 ; idpix++ ) { - ic += pathDi[idpath][idpix] ; - jc += pathDj[idpath][idpix] ; - mse.x += tex2D( tex_img_in, ic, jc ) ; - mse.y += tex2D( tex_img_in, ic, jc ) * tex2D( tex_img_in, ic, jc ) ; + ic += tex2D(tex_paths, idpix, idpath).x ; + jc += tex2D(tex_paths, idpix, idpath).y ; + zc = tex2D(tex_img_estim, jc, ic) ; + mse_data.x += zc ; + mse_data.y += zc*zc ; } // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath ) - // a ameliorer pour vitesse - mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ; - if (idpath > 0) { - if ( mse_cur < mse_min ) { - mse_min = mse_cur ; - } - } else { + // TODO cherchera ameliorer pour vitesse + mse_cur = ( mse_data.y - ( mse_data.x / lpath ) * mse_data.x ) ; + if ( mse_cur < mse_min ){ mse_min = mse_cur ; - } + val = mse_data.x ; + } } - img_out[ i*L + j ] = mse_min / lpath ; + img_out[ i*L + j ] = val / lpath ; } } - +/** + * + * \brief trace les segments sur un maillage carré + * \author zulu - AND + * + * \param[in] img_in image d'entree + * \param[in] dir tableaux des directions + * \param[in] L Largeur de l'image + * \param[in] H Hauteur de l'image + * \param[in] pas coté du maillage + * \param[in] ng niveau de gris des segments + * \param[in] r longueur des segments + * + * \param[out] img_out image + les segments superposés + * + * Kernel trivial. Ne trace rien sur les bords. + * execution sur des blocs de threads 2D et une grille 2D + * selon les dimensions de l'image + */ __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out, - unsigned int L, unsigned int H, unsigned int pas, unsigned int ng){ + unsigned int L, unsigned int H, unsigned int pas, unsigned int ng, + unsigned int r ){ // coordonnes du point dans l'image unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;; // nb de points par chemin - int lpath = PSIZE_J ; + int lpath = r ; unsigned int ic, jc, idpix ; unsigned int idpath ; @@ -306,8 +518,8 @@ __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir idpath = dir[ic*L+jc] ; img_out[ ic*L+jc ] = ng ; for ( idpix=0 ; idpix < lpath-1 ; idpix++ ){ - ic += pathDi[idpath][idpix] ; - jc += pathDj[idpath][idpix] ; + ic += tex2D(tex_paths, idpix, idpath).x ; // pathDi[idpath][idpix] ; + jc += tex2D(tex_paths, idpix, idpath).y ; // pathDj[idpath][idpix] ; img_out[ ic*L + jc ] = ng ; } } diff --git a/lniv.cvp b/lniv.cvp new file mode 100644 index 0000000..741b8ce --- /dev/null +++ b/lniv.cvp @@ -0,0 +1,148 @@ + + + + smem_3staticArrays + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 10:47:10 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + nosmem_tex + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 10:50:54 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + smem2DnoCFI + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 11:10:33 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + smemCFI_1linearArray + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 11:16:24 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + smem_zc_CFI_linearArray + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 11:30:11 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + Session19 + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 14:10:47 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + Session21 + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 14:27:50 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + Session22 + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 14:34:00 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + diff --git a/lniv_Session19_Context_0.csv b/lniv_Session19_Context_0.csv new file mode 100644 index 0000000..dedd7fe --- /dev/null +++ b/lniv_Session19_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef914083a0 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122940adae6a6780,memcpyHtoA,177.856,331,,,,,,,,,,,0,0,1048576,0 +122940adae70e2c0,_Z19kernel_calcul_pathsP4int2j,43.296,12,0.031,1,1,13908848,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122940adae719640,memcpyDtoA,4.288,6,,,,,,,,,,,0,0,576,0 +122940adae71e360,_Z23kernel_neutre_img2estimPjjj,450.88,8,0.5,64,64,0,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122940adae7969a0,memcpyAtoD,50.592,6,,,,,,,,,,,0,0,1048576,0 +122940adae7a3980,memcpyDtoA,47.392,4,,,,,,,,,,,0,0,1048576,0 +122940adae7b4960,_Z29kernel_levelines_texture_smemPjjjj,2129.12,9,0.5,64,64,6881395,8,8,1,0,308,18,0,-1,,0,47606,11,731991,67080,410,0,0,0,0,0,12802,0,0,0,272,1173117,4335 +122940adae9bd0a0,memcpyDtoA,47.712,4,,,,,,,,,,,0,0,1048576,0 +122940adae9cc6c0,_Z30kernel_estim_next_step_texturePjjjj,472.384,5,0.5,64,64,7536757,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122940adaea40640,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +122940adaea4dca0,_Z29kernel_levelines_texture_smemPjjjj,2109.98,4,0.5,64,64,6357104,8,8,1,0,308,18,0,-1,,0,47606,12,731991,66886,410,0,0,0,0,0,12804,0,0,0,272,1164819,2655 +122940adaec51900,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122940adaec5eee0,_Z30kernel_estim_next_step_texturePjjjj,472.96,4,0.5,64,64,6488175,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122940adaecd30a0,memcpyDtoA,47.744,3,,,,,,,,,,,0,0,1048576,0 +122940adaece0740,_Z29kernel_levelines_texture_smemPjjjj,2126.05,4,0.5,64,64,6619245,8,8,1,0,308,18,0,-1,,0,47606,10,731993,66947,409,0,0,0,0,0,12800,0,0,0,272,1163430,2600 +122940adaeee8260,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +122940adaeef58a0,_Z30kernel_estim_next_step_texturePjjjj,457.856,4,0.5,64,64,7602208,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122940adaef65fe0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122940adaef73620,_Z29kernel_levelines_texture_smemPjjjj,2112.03,3,0.5,64,64,6357108,8,8,1,0,308,18,0,-1,,0,47606,12,731993,66824,409,0,0,0,0,0,12788,0,0,0,272,1163358,2672 +122940adaf177b20,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122940adaf185160,_Z30kernel_estim_next_step_texturePjjjj,474.08,4,0.5,64,64,6750313,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122940adaf1f97a0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122940adaf206da0,_Z29kernel_levelines_texture_smemPjjjj,2102.98,4,0.5,64,64,6357090,8,8,1,0,308,18,0,-1,,0,47606,10,731993,66846,410,0,0,0,0,0,12815,0,0,0,272,1165192,2643 +122940adaf408ec0,memcpyDtoA,47.808,4,,,,,,,,,,,0,0,1048576,0 +122940adaf4165a0,_Z30kernel_estim_next_step_texturePjjjj,479.456,4,0.5,64,64,7602276,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122940adaf48c0e0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122940adaf499720,_Z29kernel_levelines_texture_smemPjjjj,2106.5,3,0.5,64,64,7798900,8,8,1,0,308,18,0,-1,,0,47434,12,729319,66968,410,0,0,0,0,0,12802,0,0,0,271,1163817,2749 +122940adaf69c5e0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122940adaf6a9c20,_Z30kernel_estim_next_step_texturePjjjj,484.128,4,0.5,64,64,7274600,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122940adaf720980,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122940adaf72df80,_Z29kernel_levelines_texture_smemPjjjj,2115.81,3,0.5,64,64,6553710,8,8,1,0,308,18,0,-1,,0,47434,10,729328,66761,410,0,0,0,0,0,12804,0,0,0,271,1164886,2588 +122940adaf9332c0,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122940adaf9408c0,_Z30kernel_estim_next_step_texturePjjjj,481.312,4,0.5,64,64,6488169,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adaf9b6b40,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122940adaf9c41a0,_Z29kernel_levelines_texture_smemPjjjj,2109.15,4,0.5,64,64,7274605,8,8,1,0,308,18,0,-1,,0,47256,10,726617,66445,409,0,0,0,0,0,12800,0,0,0,270,1163404,2626 +122940adafbc7ac0,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +122940adafbd5120,_Z30kernel_estim_next_step_texturePjjjj,477.376,4,0.5,64,64,7733327,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adafc4a420,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122940adafc57a20,_Z29kernel_levelines_texture_smemPjjjj,2113.09,4,0.5,64,64,2097253,8,8,1,0,308,18,0,-1,,0,47256,10,726618,66738,409,0,0,0,0,0,12788,0,0,0,270,1163359,2671 +122940adafe5c2e0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adafe698c0,_Z30kernel_estim_next_step_texturePjjjj,476.672,3,0.5,64,64,7209065,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adafede940,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122940adafeebf60,_Z29kernel_levelines_texture_smemPjjjj,2110.3,4,0.5,64,64,6619245,8,8,1,0,308,18,0,-1,,0,47256,10,726610,66368,410,0,0,0,0,0,12815,0,0,0,270,1165183,2652 +122940adb00efd00,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122940adb00fd300,_Z30kernel_estim_next_step_texturePjjjj,468.864,4,0.5,64,64,7536672,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122940adb01704c0,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122940adb017dae0,_Z29kernel_levelines_texture_smemPjjjj,2115.36,4,0.5,64,64,2097252,8,8,1,0,308,18,0,-1,,0,47256,10,726611,66084,410,0,0,0,0,0,12802,0,0,0,270,1163872,2694 +122940adb0382c40,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122940adb0390220,_Z30kernel_estim_next_step_texturePjjjj,473.44,4,0.5,64,64,7274614,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122940adb04045c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122940adb0411bc0,_Z29kernel_levelines_texture_smemPjjjj,2116.22,4,0.5,64,64,6357024,8,8,1,0,308,18,0,-1,,0,47256,10,726610,66314,410,0,0,0,0,0,12804,0,0,0,270,1164898,2576 +122940adb0617080,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adb0624680,_Z30kernel_estim_next_step_texturePjjjj,466.592,4,0.5,64,64,7143456,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adb0696f60,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122940adb06a45a0,_Z29kernel_levelines_texture_smemPjjjj,2121.09,4,0.5,64,64,6553701,8,8,1,0,308,18,0,-1,,0,47256,10,726613,66552,409,0,0,0,0,0,12800,0,0,0,270,1163359,2671 +122940adb08aad80,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122940adb08b83a0,_Z30kernel_estim_next_step_texturePjjjj,465.152,4,0.5,64,64,7733359,8,8,1,0,36,9,0,-1,,0,822,0,15627,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adb092a700,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adb0937d40,_Z29kernel_levelines_texture_smemPjjjj,2106.34,4,0.5,64,64,2097260,8,8,1,0,308,18,0,-1,,0,47256,10,726614,66553,409,0,0,0,0,0,12788,0,0,0,270,1163415,2615 +122940adb0b3ab60,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adb0b48140,_Z30kernel_estim_next_step_texturePjjjj,476.896,4,0.5,64,64,7143525,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122940adb0bbd2a0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adb0bca8a0,_Z29kernel_levelines_texture_smemPjjjj,2109.22,4,0.5,64,64,7274598,8,8,1,0,308,18,0,-1,,0,47606,13,731993,67183,410,0,0,0,0,0,12815,0,0,0,272,1165195,2640 +122940adb0dce280,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122940adb0ddb880,_Z30kernel_estim_next_step_texturePjjjj,483.456,3,0.5,64,64,6619235,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122940adb0e8f400,memcpyAtoH,11.52,29,,,,,,,,,,,0,0,576,0 +122940adb0e98720,memcpyDtoH,191.68,736,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_Session21_Context_0.csv b/lniv_Session21_Context_0.csv new file mode 100644 index 0000000..5f60f58 --- /dev/null +++ b/lniv_Session21_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef91ef4898 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +1229419bcf456a00,memcpyHtoA,177.728,334,,,,,,,,,,,0,0,1048576,0 +1229419bcf4bf660,_Z19kernel_calcul_pathsP4int2j,43.392,12,0.031,1,1,12973520,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +1229419bcf4caa20,memcpyDtoA,4.32,6,,,,,,,,,,,0,0,576,0 +1229419bcf4cf760,_Z23kernel_neutre_img2estimPjjj,440.768,9,0.5,64,64,-2147483648,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +1229419bcf544f00,memcpyAtoD,50.4,6,,,,,,,,,,,0,0,1048576,0 +1229419bcf551e20,memcpyDtoA,47.52,3,,,,,,,,,,,0,0,1048576,0 +1229419bcf562e80,_Z29kernel_levelines_texture_smemPjjjj,2074.66,9,0.5,64,64,8916960,8,8,1,0,36,18,0,-1,,0,47606,11,703977,0,410,0,0,0,0,0,12802,0,0,0,272,1173148,4304 +1229419bcf75e0e0,memcpyDtoA,47.744,4,,,,,,,,,,,0,0,1048576,0 +1229419bcf76d700,_Z30kernel_estim_next_step_texturePjjjj,475.36,6,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +1229419bcf7e2220,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +1229419bcf7ef860,_Z29kernel_levelines_texture_smemPjjjj,2052.22,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,12,703981,0,410,0,0,0,0,0,12804,0,0,0,272,1164805,2669 +1229419bcf9e5320,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +1229419bcf9f2900,_Z30kernel_estim_next_step_texturePjjjj,457.632,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +1229419bcfa62f40,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bcfa705a0,_Z29kernel_levelines_texture_smemPjjjj,2059.81,4,0.5,64,64,8466560,8,8,1,0,36,18,0,-1,,0,47606,10,703983,0,409,0,0,0,0,0,12800,0,0,0,272,1163425,2605 +1229419bcfc67e00,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bcfc753e0,_Z30kernel_estim_next_step_texturePjjjj,484.128,3,0.5,64,64,10922,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +1229419bcfcec140,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bcfcf9720,_Z29kernel_levelines_texture_smemPjjjj,2054.18,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,12,703989,0,409,0,0,0,0,0,12788,0,0,0,272,1163336,2694 +1229419bcfeef9a0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +1229419bcfefcfe0,_Z30kernel_estim_next_step_texturePjjjj,469.376,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +1229419bcff703a0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bcff7d9a0,_Z29kernel_levelines_texture_smemPjjjj,2055.1,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,10,703984,0,410,0,0,0,0,0,12815,0,0,0,272,1165156,2679 +1229419bd0173fa0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bd0181600,_Z30kernel_estim_next_step_texturePjjjj,477.216,9,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +1229419bd01f6840,memcpyDtoA,47.712,4,,,,,,,,,,,0,0,1048576,0 +1229419bd0203ec0,_Z29kernel_levelines_texture_smemPjjjj,2056.77,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47434,12,701415,0,410,0,0,0,0,0,12802,0,0,0,271,1163891,2675 +1229419bd03fab40,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +1229419bd04081a0,_Z30kernel_estim_next_step_texturePjjjj,484.992,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +1229419bd047f280,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +1229419bd048c8e0,_Z29kernel_levelines_texture_smemPjjjj,2049.92,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47434,10,701416,0,410,0,0,0,0,0,12804,0,0,0,271,1164883,2591 +1229419bd0681ac0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bd068f0e0,_Z30kernel_estim_next_step_texturePjjjj,486.752,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd07068a0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +1229419bd0713e80,_Z29kernel_levelines_texture_smemPjjjj,2068.54,4,0.5,64,64,7143535,8,8,1,0,36,18,0,-1,,0,47256,10,698805,0,409,0,0,0,0,0,12800,0,0,0,270,1163428,2602 +1229419bd090d900,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +1229419bd091af60,_Z30kernel_estim_next_step_texturePjjjj,489.152,3,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd09930c0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +1229419bd09a0700,_Z29kernel_levelines_texture_smemPjjjj,2058.94,3,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47256,10,698807,0,409,0,0,0,0,0,12788,0,0,0,270,1163396,2634 +1229419bd0b97c40,memcpyDtoA,47.52,3,,,,,,,,,,,0,0,1048576,0 +1229419bd0ba5220,_Z30kernel_estim_next_step_texturePjjjj,479.776,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd0c1aea0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bd0c284a0,_Z29kernel_levelines_texture_smemPjjjj,2055.78,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47256,10,698809,0,410,0,0,0,0,0,12815,0,0,0,270,1165209,2626 +1229419bd0e1ed60,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +1229419bd0e2c3e0,_Z30kernel_estim_next_step_texturePjjjj,471.392,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +1229419bd0e9ff80,memcpyDtoA,47.808,3,,,,,,,,,,,0,0,1048576,0 +1229419bd0ead6a0,_Z29kernel_levelines_texture_smemPjjjj,2047.3,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47256,10,698801,0,410,0,0,0,0,0,12802,0,0,0,270,1163935,2631 +1229419bd10a1e20,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +1229419bd10af420,_Z30kernel_estim_next_step_texturePjjjj,476.864,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15627,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +1229419bd1124540,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +1229419bd1131ba0,_Z29kernel_levelines_texture_smemPjjjj,2057.82,4,0.5,64,64,-1631,8,8,1,0,36,18,0,-1,,0,47256,10,698807,0,410,0,0,0,0,0,12804,0,0,0,270,1164939,2535 +1229419bd1328c60,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +1229419bd1336240,_Z30kernel_estim_next_step_texturePjjjj,475.936,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd13aafa0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +1229419bd13b85a0,_Z29kernel_levelines_texture_smemPjjjj,2057.95,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47256,10,698806,0,409,0,0,0,0,0,12800,0,0,0,270,1163422,2608 +1229419bd15af6e0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +1229419bd15bcd60,_Z30kernel_estim_next_step_texturePjjjj,493.888,4,0.5,64,64,856,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd16360e0,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +1229419bd1643760,_Z29kernel_levelines_texture_smemPjjjj,2053.98,3,0.5,64,64,12758580,8,8,1,0,36,18,0,-1,,0,47256,10,698810,0,409,0,0,0,0,0,12788,0,0,0,270,1163489,2541 +1229419bd1839940,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +1229419bd1846f40,_Z30kernel_estim_next_step_texturePjjjj,474.208,3,0.5,64,64,32767,8,8,1,0,36,9,0,-1,,0,822,0,15628,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229419bd18bb600,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +1229419bd18c8c40,_Z29kernel_levelines_texture_smemPjjjj,2065.76,4,0.5,64,64,2,8,8,1,0,36,18,0,-1,,0,47606,13,703985,0,410,0,0,0,0,0,12815,0,0,0,272,1165162,2673 +1229419bd1ac1be0,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +1229419bd1acf1e0,_Z30kernel_estim_next_step_texturePjjjj,465.952,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +1229419bd1b7e360,memcpyAtoH,11.52,30,,,,,,,,,,,0,0,576,0 +1229419bd1b87d00,memcpyDtoH,191.232,736,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_Session22_Context_0.csv b/lniv_Session22_Context_0.csv new file mode 100644 index 0000000..fe8eba7 --- /dev/null +++ b/lniv_Session22_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef92065818 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122941f213dd9800,memcpyHtoA,177.792,336,,,,,,,,,,,0,0,1048576,0 +122941f213e40de0,_Z19kernel_calcul_pathsP4int2j,43.2,12,0.031,1,1,-1386402264,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122941f213e4c140,memcpyDtoA,4.32,6,,,,,,,,,,,0,0,576,0 +122941f213e50ea0,_Z23kernel_neutre_img2estimPjjj,446.944,9,0.5,64,64,0,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122941f213ec8740,memcpyAtoD,50.496,6,,,,,,,,,,,0,0,1048576,0 +122941f213ed56c0,memcpyDtoA,47.456,4,,,,,,,,,,,0,0,1048576,0 +122941f213ee6700,_Z29kernel_levelines_texture_smemPjjjj,2749.79,9,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,11,743152,187680,410,0,0,0,0,0,12802,0,0,0,272,1571115,7958 +122941f2141866a0,memcpyDtoA,47.84,4,,,,,,,,,,,0,0,1048576,0 +122941f214195d60,_Z30kernel_estim_next_step_texturePjjjj,472.736,6,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122941f214209e60,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f214217440,_Z29kernel_levelines_texture_smemPjjjj,2289.44,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,12,743153,152442,410,0,0,0,0,0,12804,0,0,0,272,1256596,17308 +122941f214446dc0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122941f214454400,_Z30kernel_estim_next_step_texturePjjjj,480.32,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122941f2144ca2a0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122941f2144d78c0,_Z29kernel_levelines_texture_smemPjjjj,2174.62,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,10,743148,144198,409,0,0,0,0,0,12800,0,0,0,272,1174004,1661 +122941f2146eb1a0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f2146f8780,_Z30kernel_estim_next_step_texturePjjjj,478.304,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122941f21476de20,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f21477b420,_Z29kernel_levelines_texture_smemPjjjj,2169.34,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,12,743151,144449,409,0,0,0,0,0,12788,0,0,0,272,1172650,1348 +122941f21498d8c0,memcpyDtoA,47.52,3,,,,,,,,,,,0,0,1048576,0 +122941f21499ae80,_Z30kernel_estim_next_step_texturePjjjj,482.432,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122941f214a11560,memcpyDtoA,47.488,3,,,,,,,,,,,0,0,1048576,0 +122941f214a1eb00,_Z29kernel_levelines_texture_smemPjjjj,2166.24,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,10,743150,144061,410,0,0,0,0,0,12815,0,0,0,272,1171928,1205 +122941f214c30340,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122941f214c3d900,_Z30kernel_estim_next_step_texturePjjjj,484.064,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122941f214cb4640,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122941f214cc1ca0,_Z29kernel_levelines_texture_smemPjjjj,2165.44,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47434,12,740447,143891,410,0,0,0,0,0,12802,0,0,0,271,1169458,1131 +122941f214ed31c0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122941f214ee0840,_Z30kernel_estim_next_step_texturePjjjj,475.104,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122941f214f55260,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f214f62840,_Z29kernel_levelines_texture_smemPjjjj,2169.92,3,0.5,64,64,19,8,8,1,0,292,20,0,-1,,0,47434,10,740437,144692,410,0,0,0,0,0,12804,0,0,0,271,1171196,1166 +122941f215174f40,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122941f215182520,_Z30kernel_estim_next_step_texturePjjjj,481.824,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f2151f8980,memcpyDtoA,47.744,3,,,,,,,,,,,0,0,1048576,0 +122941f215206020,_Z29kernel_levelines_texture_smemPjjjj,2174.78,3,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737697,142626,409,0,0,0,0,0,12800,0,0,0,270,1171005,1147 +122941f2154199c0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122941f215427060,_Z30kernel_estim_next_step_texturePjjjj,473.44,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f21549b400,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f2154a89e0,_Z29kernel_levelines_texture_smemPjjjj,2172.13,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737693,142582,409,0,0,0,0,0,12788,0,0,0,270,1170316,1122 +122941f2156bb900,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f2156c8f00,_Z30kernel_estim_next_step_texturePjjjj,478.72,4,0.5,64,64,512,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f21573e740,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122941f21574bda0,_Z29kernel_levelines_texture_smemPjjjj,2175.81,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737683,142961,410,0,0,0,0,0,12815,0,0,0,270,1172017,1158 +122941f21595fb20,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122941f21596d120,_Z30kernel_estim_next_step_texturePjjjj,484.352,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122941f2159e3f80,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122941f2159f1580,_Z29kernel_levelines_texture_smemPjjjj,2171.97,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737684,143028,410,0,0,0,0,0,12802,0,0,0,270,1170115,1118 +122941f215c04400,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122941f215c11a40,_Z30kernel_estim_next_step_texturePjjjj,473.408,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15627,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122941f215c85dc0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122941f215c93420,_Z29kernel_levelines_texture_smemPjjjj,2178.5,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737687,143703,410,0,0,0,0,0,12804,0,0,0,270,1171501,1189 +122941f215ea7c20,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122941f215eb5200,_Z30kernel_estim_next_step_texturePjjjj,475.584,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f215f29e00,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +122941f215f37460,_Z29kernel_levelines_texture_smemPjjjj,2162.05,3,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737680,142311,409,0,0,0,0,0,12800,0,0,0,270,1170773,1141 +122941f216147c20,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122941f216155220,_Z30kernel_estim_next_step_texturePjjjj,473.152,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f2161c94a0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122941f2161d6a80,_Z29kernel_levelines_texture_smemPjjjj,2169.7,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47256,10,737687,142117,409,0,0,0,0,0,12788,0,0,0,270,1169968,1132 +122941f2163e9020,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122941f2163f66c0,_Z30kernel_estim_next_step_texturePjjjj,491.52,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122941f21646f140,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122941f21647c740,_Z29kernel_levelines_texture_smemPjjjj,2164.67,4,0.5,64,64,0,8,8,1,0,292,20,0,-1,,0,47606,13,743152,142819,410,0,0,0,0,0,12815,0,0,0,272,1171022,1143 +122941f21668d940,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122941f21669af80,_Z30kernel_estim_next_step_texturePjjjj,486.816,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122941f21674db40,memcpyAtoH,11.552,30,,,,,,,,,,,0,0,576,0 +122941f216756cc0,memcpyDtoH,191.712,745,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_nosmem_tex_Context_0.csv b/lniv_nosmem_tex_Context_0.csv new file mode 100644 index 0000000..733b9a7 --- /dev/null +++ b/lniv_nosmem_tex_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef88fd8ba8 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122935c55701cc60,memcpyHtoA,177.568,331,,,,,,,,,,,0,0,1048576,0 +122935c557085020,_Z19kernel_calcul_pathsP4int2j,43.168,13,0.031,1,1,11485616,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122935c557090340,memcpyDtoA,4.32,5,,,,,,,,,,,0,0,576,0 +122935c5570950c0,_Z23kernel_neutre_img2estimPjjj,441.44,9,0.5,64,64,1,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122935c55710b3a0,memcpyAtoD,50.464,6,,,,,,,,,,,0,0,1048576,0 +122935c557118300,memcpyDtoA,47.392,4,,,,,,,,,,,0,0,1048576,0 +122935c557129320,_Z24kernel_levelines_texturePjjjj,2217.66,10,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,11,748312,0,410,0,0,0,0,0,12802,0,0,0,272,1198877,41827 +122935c557347440,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122935c557356a00,_Z30kernel_estim_next_step_texturePjjjj,463.072,6,0.5,64,64,13,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935c5573c8540,memcpyDtoA,47.52,3,,,,,,,,,,,0,0,1048576,0 +122935c5573d5b20,_Z24kernel_levelines_texturePjjjj,2217.76,4,0.5,64,64,11833200,8,8,1,0,36,18,0,-1,,0,47606,12,748315,0,410,0,0,0,0,0,12804,0,0,0,272,1198392,43464 +122935c5575f3ca0,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935c557601280,_Z30kernel_estim_next_step_texturePjjjj,471.84,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935c557674fe0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c5576825e0,_Z24kernel_levelines_texturePjjjj,2230.08,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,10,748312,0,409,0,0,0,0,0,12800,0,0,0,272,1195784,44536 +122935c5578a37e0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c5578b0de0,_Z30kernel_estim_next_step_texturePjjjj,466.176,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935c557923540,memcpyDtoA,47.712,4,,,,,,,,,,,0,0,1048576,0 +122935c557930bc0,_Z24kernel_levelines_texturePjjjj,2213.15,4,0.5,64,64,0,8,8,1,0,36,18,0,-1,,0,47606,12,748313,0,409,0,0,0,0,0,12788,0,0,0,272,1198714,41606 +122935c557b4dba0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935c557b5b1c0,_Z30kernel_estim_next_step_texturePjjjj,471.072,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935c557bcec20,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c557bdc260,_Z24kernel_levelines_texturePjjjj,2206.82,4,0.5,64,64,11460208,8,8,1,0,36,18,0,-1,,0,47606,10,748330,0,410,0,0,0,0,0,12815,0,0,0,272,1198656,43584 +122935c557df7980,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122935c557e04fc0,_Z30kernel_estim_next_step_texturePjjjj,478.304,3,0.5,64,64,10922,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935c557e7a6c0,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122935c557e87d40,_Z24kernel_levelines_texturePjjjj,2228.32,4,0.5,64,64,32,8,8,1,0,36,18,0,-1,,0,47434,12,745579,0,410,0,0,0,0,0,12802,0,0,0,271,1198277,42427 +122935c5580a87e0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122935c5580b5de0,_Z30kernel_estim_next_step_texturePjjjj,467.04,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935c5581288e0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122935c558135ea0,_Z24kernel_levelines_texturePjjjj,2214.69,4,0.5,64,64,-1053464,8,8,1,0,36,18,0,-1,,0,47434,10,745588,0,410,0,0,0,0,0,12804,0,0,0,271,1199501,42355 +122935c558353400,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122935c558360a40,_Z30kernel_estim_next_step_texturePjjjj,491.712,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c5583d9560,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935c5583e6b60,_Z24kernel_levelines_texturePjjjj,2220.61,4,0.5,64,64,11831600,8,8,1,0,36,18,0,-1,,0,47256,10,742810,0,409,0,0,0,0,0,12800,0,0,0,270,1195661,44659 +122935c558605840,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122935c558612e80,_Z30kernel_estim_next_step_texturePjjjj,472.992,3,0.5,64,64,256,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c558687060,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c558694680,_Z24kernel_levelines_texturePjjjj,2218.37,4,0.5,64,64,8373664,8,8,1,0,36,18,0,-1,,0,47256,10,742813,0,409,0,0,0,0,0,12788,0,0,0,270,1197654,42666 +122935c5588b2a60,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c5588c0060,_Z30kernel_estim_next_step_texturePjjjj,478.496,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c5589357c0,memcpyDtoA,47.52,4,,,,,,,,,,,0,0,1048576,0 +122935c558942d80,_Z24kernel_levelines_texturePjjjj,2210.37,3,0.5,64,64,144,8,8,1,0,36,18,0,-1,,0,47256,10,742811,0,410,0,0,0,0,0,12815,0,0,0,270,1198542,43698 +122935c558b5f200,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935c558b6c800,_Z30kernel_estim_next_step_texturePjjjj,467.968,4,0.5,64,64,48,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122935c558bdf660,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +122935c558beccc0,_Z24kernel_levelines_texturePjjjj,2221.98,4,0.5,64,64,11565802,8,8,1,0,36,18,0,-1,,0,47256,10,742809,0,410,0,0,0,0,0,12802,0,0,0,270,1198356,42348 +122935c558e0bea0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935c558e194a0,_Z30kernel_estim_next_step_texturePjjjj,468.512,4,0.5,64,64,7602289,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122935c558e8c500,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935c558e99b80,_Z24kernel_levelines_texturePjjjj,2216.1,4,0.5,64,64,6226030,8,8,1,0,36,18,0,-1,,0,47256,10,742809,0,410,0,0,0,0,0,12804,0,0,0,270,1199237,42619 +122935c5590b76e0,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935c5590c4cc0,_Z30kernel_estim_next_step_texturePjjjj,477.856,4,0.5,64,64,3211314,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c55913a1c0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122935c5591477a0,_Z24kernel_levelines_texturePjjjj,2213.6,4,0.5,64,64,3670069,8,8,1,0,36,18,0,-1,,0,47256,10,742809,0,409,0,0,0,0,0,12800,0,0,0,270,1196806,43514 +122935c5593648e0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122935c559371f40,_Z30kernel_estim_next_step_texturePjjjj,484.352,4,0.5,64,64,3145780,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c5593e8d80,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122935c5593f6400,_Z24kernel_levelines_texturePjjjj,2214.46,4,0.5,64,64,6225971,8,8,1,0,36,18,0,-1,,0,47256,10,742810,0,409,0,0,0,0,0,12788,0,0,0,270,1197263,43057 +122935c5596138a0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122935c559620f20,_Z30kernel_estim_next_step_texturePjjjj,485.952,4,0.5,64,64,3145773,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935c5596983a0,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +122935c5596a5a40,_Z24kernel_levelines_texturePjjjj,2219.36,4,0.5,64,64,-767361184,8,8,1,0,36,18,0,-1,,0,47606,13,748312,0,410,0,0,0,0,0,12815,0,0,0,272,1199612,42628 +122935c5598c41e0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935c5598d1860,_Z30kernel_estim_next_step_texturePjjjj,469.504,4,0.5,64,64,48,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935c559988ea0,memcpyAtoH,11.584,33,,,,,,,,,,,0,0,576,0 +122935c5599936c0,memcpyDtoH,191.04,741,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_smem2DnoCFI_Context_0.csv b/lniv_smem2DnoCFI_Context_0.csv new file mode 100644 index 0000000..764af7e --- /dev/null +++ b/lniv_smem2DnoCFI_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef899f0960 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122936d7f2be0c20,memcpyHtoA,177.248,331,,,,,,,,,,,0,0,1048576,0 +122936d7f2c48180,_Z19kernel_calcul_pathsP4int2j,43.2,12,0.031,1,1,12468336,1,1,1,0,28,11,0,-1,,0,168,0,1588,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122936d7f2c53480,memcpyDtoA,4.256,6,,,,,,,,,,,0,0,576,0 +122936d7f2c581a0,_Z23kernel_neutre_img2estimPjjj,432.8,8,0.5,64,64,256,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122936d7f2ccc5a0,memcpyAtoD,50.4,6,,,,,,,,,,,0,0,1048576,0 +122936d7f2cd94c0,memcpyDtoA,47.488,4,,,,,,,,,,,0,0,1048576,0 +122936d7f2cea560,_Z29kernel_levelines_texture_smemPjjjj,3264.77,9,0.5,64,64,11454096,8,8,1,0,552,18,0,-1,,0,47606,11,788295,862664,410,0,0,0,0,0,12802,0,0,0,272,1224934,31164 +122936d7f30080a0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122936d7f30176c0,_Z30kernel_estim_next_step_texturePjjjj,483.904,6,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122936d7f308e340,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122936d7f309b960,_Z29kernel_levelines_texture_smemPjjjj,3259.26,4,0.5,64,64,64,8,8,1,0,552,18,0,-1,,0,47606,12,788295,858288,410,0,0,0,0,0,12804,0,0,0,272,1220536,26811 +122936d7f33b7f20,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f33c5500,_Z30kernel_estim_next_step_texturePjjjj,485.856,5,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122936d7f343c940,memcpyDtoA,47.776,4,,,,,,,,,,,0,0,1048576,0 +122936d7f344a000,_Z29kernel_levelines_texture_smemPjjjj,3269.73,3,0.5,64,64,0,8,8,1,0,552,18,0,-1,,0,47606,10,788295,857612,409,0,0,0,0,0,12800,0,0,0,272,1214814,27588 +122936d7f3768ea0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122936d7f37764c0,_Z30kernel_estim_next_step_texturePjjjj,481.024,4,0.5,64,64,31,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122936d7f37ec680,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f37f9c60,_Z29kernel_levelines_texture_smemPjjjj,3263.46,3,0.5,64,64,147,8,8,1,0,552,18,0,-1,,0,47606,12,788302,855382,409,0,0,0,0,0,12788,0,0,0,272,1215865,24785 +122936d7f3b172a0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122936d7f3b248e0,_Z30kernel_estim_next_step_texturePjjjj,470.016,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122936d7f3b97f60,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122936d7f3ba55a0,_Z29kernel_levelines_texture_smemPjjjj,3258.18,4,0.5,64,64,12305584,8,8,1,0,552,18,0,-1,,0,47606,10,788297,859076,410,0,0,0,0,0,12815,0,0,0,272,1215402,27600 +122936d7f3ec1740,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122936d7f3eced80,_Z30kernel_estim_next_step_texturePjjjj,465.536,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122936d7f3f41260,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f3f4e860,_Z29kernel_levelines_texture_smemPjjjj,3266.59,4,0.5,64,64,117,8,8,1,0,552,18,0,-1,,0,47434,12,785416,855772,410,0,0,0,0,0,12802,0,0,0,271,1216005,25733 +122936d7f426cb40,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f427a120,_Z30kernel_estim_next_step_texturePjjjj,478.304,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122936d7f42ef7c0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122936d7f42fce60,_Z29kernel_levelines_texture_smemPjjjj,3274.46,3,0.5,64,64,192,8,8,1,0,552,18,0,-1,,0,47434,10,785418,856700,410,0,0,0,0,0,12804,0,0,0,271,1218747,24325 +122936d7f461cfa0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f462a600,_Z30kernel_estim_next_step_texturePjjjj,476.576,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f469f600,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f46acc20,_Z29kernel_levelines_texture_smemPjjjj,3270.85,3,0.5,64,64,-1404411488,8,8,1,0,552,18,0,-1,,0,47256,10,782499,851609,409,0,0,0,0,0,12800,0,0,0,270,1215219,26327 +122936d7f49cbfa0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122936d7f49d95a0,_Z30kernel_estim_next_step_texturePjjjj,478.88,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f4a4eea0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122936d7f4a5c500,_Z29kernel_levelines_texture_smemPjjjj,3272.86,3,0.5,64,64,0,8,8,1,0,552,18,0,-1,,0,47256,10,782499,851674,409,0,0,0,0,0,12788,0,0,0,270,1215877,25735 +122936d7f4d7c000,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f4d89640,_Z30kernel_estim_next_step_texturePjjjj,479.904,4,0.5,64,64,35,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f4dff320,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122936d7f4e0c920,_Z29kernel_levelines_texture_smemPjjjj,3263.14,4,0.5,64,64,145,8,8,1,0,552,18,0,-1,,0,47256,10,782501,852832,410,0,0,0,0,0,12815,0,0,0,270,1216391,27119 +122936d7f5129e00,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122936d7f5137460,_Z30kernel_estim_next_step_texturePjjjj,474.112,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122936d7f51abaa0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122936d7f51b90c0,_Z29kernel_levelines_texture_smemPjjjj,3255.3,3,0.5,64,64,0,8,8,1,0,552,18,0,-1,,0,47256,10,782499,851161,410,0,0,0,0,0,12802,0,0,0,270,1216487,25419 +122936d7f54d4700,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f54e1ce0,_Z30kernel_estim_next_step_texturePjjjj,478.432,3,0.5,64,64,37,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122936d7f5557420,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122936d7f5564a60,_Z29kernel_levelines_texture_smemPjjjj,3263.68,4,0.5,64,64,1102,8,8,1,0,552,18,0,-1,,0,47256,10,782499,851691,410,0,0,0,0,0,12804,0,0,0,270,1215992,27116 +122936d7f5882180,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122936d7f588f760,_Z30kernel_estim_next_step_texturePjjjj,470.944,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f5903140,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122936d7f5910760,_Z29kernel_levelines_texture_smemPjjjj,3257.12,3,0.5,64,64,9661136,8,8,1,0,552,18,0,-1,,0,47256,10,782499,851971,409,0,0,0,0,0,12800,0,0,0,270,1216412,25128 +122936d7f5c2c4c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f5c39ac0,_Z30kernel_estim_next_step_texturePjjjj,483.392,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15627,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f5cb0560,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f5cbdb60,_Z29kernel_levelines_texture_smemPjjjj,3269.76,4,0.5,64,64,12157856,8,8,1,0,552,18,0,-1,,0,47256,10,782500,855086,409,0,0,0,0,0,12788,0,0,0,270,1217387,24205 +122936d7f5fdca40,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122936d7f5fea060,_Z30kernel_estim_next_step_texturePjjjj,466.784,3,0.5,64,64,32767,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122936d7f605ca20,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122936d7f606a020,_Z29kernel_levelines_texture_smemPjjjj,3273.89,5,0.5,64,64,9749904,8,8,1,0,552,18,0,-1,,0,47606,13,788296,857119,410,0,0,0,0,0,12815,0,0,0,272,1216502,27016 +122936d7f6389f20,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122936d7f6397580,_Z30kernel_estim_next_step_texturePjjjj,478.08,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122936d7f6447fa0,memcpyAtoH,11.616,29,,,,,,,,,,,0,0,576,0 +122936d7f6451160,memcpyDtoH,191.392,744,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_smemCFI_1linearArray_Context_0.csv b/lniv_smemCFI_1linearArray_Context_0.csv new file mode 100644 index 0000000..23169a9 --- /dev/null +++ b/lniv_smemCFI_1linearArray_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef89bda598 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122937299dd13280,memcpyHtoA,177.568,333,,,,,,,,,,,0,0,1048576,0 +122937299dd7b040,_Z19kernel_calcul_pathsP4int2j,42.976,12,0.031,1,1,12201664,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122937299dd86260,memcpyDtoA,4.288,6,,,,,,,,,,,0,0,576,0 +122937299dd8afa0,_Z23kernel_neutre_img2estimPjjj,444.608,8,0.5,64,64,1074266112,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122937299de016c0,memcpyAtoD,50.592,6,,,,,,,,,,,0,0,1048576,0 +122937299de0e6e0,memcpyDtoA,47.424,4,,,,,,,,,,,0,0,1048576,0 +122937299de1f760,_Z29kernel_levelines_texture_smemPjjjj,2694.56,10,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47606,11,789386,391913,410,0,0,0,0,0,12802,0,0,0,272,1223244,32854 +122937299e0b1f60,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122937299e0c1560,_Z30kernel_estim_next_step_texturePjjjj,471.616,6,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122937299e135200,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937299e142860,_Z29kernel_levelines_texture_smemPjjjj,2685.31,4,0.5,64,64,1,8,8,1,0,584,18,0,-1,,0,47606,12,789386,390037,410,0,0,0,0,0,12804,0,0,0,272,1219847,27708 +122937299e3d2c20,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122937299e3e0260,_Z30kernel_estim_next_step_texturePjjjj,486.784,4,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122937299e457a40,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937299e465040,_Z29kernel_levelines_texture_smemPjjjj,2667.74,4,0.5,64,64,1,8,8,1,0,584,18,0,-1,,0,47606,10,789383,389311,409,0,0,0,0,0,12800,0,0,0,272,1213376,27080 +122937299e6f0fa0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937299e6fe600,_Z30kernel_estim_next_step_texturePjjjj,470.112,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122937299e771cc0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122937299e77f2a0,_Z29kernel_levelines_texture_smemPjjjj,2670.72,3,0.5,64,64,1,8,8,1,0,584,18,0,-1,,0,47606,12,789386,388524,409,0,0,0,0,0,12788,0,0,0,272,1215272,26206 +122937299ea0bd60,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122937299ea19340,_Z30kernel_estim_next_step_texturePjjjj,480.096,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122937299ea8f0e0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122937299ea9c740,_Z29kernel_levelines_texture_smemPjjjj,2671.14,4,0.5,64,64,1,8,8,1,0,584,18,0,-1,,0,47606,10,789388,389600,410,0,0,0,0,0,12815,0,0,0,272,1218236,25318 +122937299ed293c0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122937299ed36a20,_Z30kernel_estim_next_step_texturePjjjj,476.704,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122937299edabac0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122937299edb90a0,_Z29kernel_levelines_texture_smemPjjjj,2671.3,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47434,12,786502,388436,410,0,0,0,0,0,12802,0,0,0,271,1213625,28427 +122937299f045da0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937299f0533c0,_Z30kernel_estim_next_step_texturePjjjj,472.064,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122937299f0c7200,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122937299f0d4860,_Z29kernel_levelines_texture_smemPjjjj,2677.92,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47434,10,786502,388243,410,0,0,0,0,0,12804,0,0,0,271,1216963,26153 +122937299f362f40,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122937299f370540,_Z30kernel_estim_next_step_texturePjjjj,467.456,3,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937299f3e3180,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122937299f3f07e0,_Z29kernel_levelines_texture_smemPjjjj,2675.04,4,0.5,64,64,11741088,8,8,1,0,584,18,0,-1,,0,47256,10,783580,385314,409,0,0,0,0,0,12800,0,0,0,270,1216083,25553 +122937299f67e3a0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937299f68ba00,_Z30kernel_estim_next_step_texturePjjjj,475.392,3,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937299f700560,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937299f70dba0,_Z29kernel_levelines_texture_smemPjjjj,2662.05,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47256,10,783580,384899,409,0,0,0,0,0,12788,0,0,0,270,1214437,27265 +122937299f998480,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937299f9a5aa0,_Z30kernel_estim_next_step_texturePjjjj,477.12,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937299fa1acc0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937299fa282c0,_Z29kernel_levelines_texture_smemPjjjj,2673.63,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47256,10,783582,384740,410,0,0,0,0,0,12815,0,0,0,270,1216791,26747 +122937299fcb58e0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122937299fcc2ee0,_Z30kernel_estim_next_step_texturePjjjj,481.792,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122937299fd39340,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122937299fd46940,_Z29kernel_levelines_texture_smemPjjjj,2675.68,3,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47256,10,783579,383647,410,0,0,0,0,0,12802,0,0,0,270,1215897,26147 +122937299ffd47c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937299ffe1dc0,_Z30kernel_estim_next_step_texturePjjjj,486.4,3,0.5,64,64,10922,8,8,1,0,36,9,0,-1,,0,822,0,15627,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +12293729a0059400,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +12293729a00669e0,_Z29kernel_levelines_texture_smemPjjjj,2668.42,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47256,10,783580,384611,410,0,0,0,0,0,12804,0,0,0,270,1216801,26315 +12293729a02f2ba0,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +12293729a0300220,_Z30kernel_estim_next_step_texturePjjjj,465.92,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +12293729a0372860,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +12293729a037fea0,_Z29kernel_levelines_texture_smemPjjjj,2667.1,3,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47256,10,783579,385331,409,0,0,0,0,0,12800,0,0,0,270,1215586,26050 +12293729a060bb80,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +12293729a06191e0,_Z30kernel_estim_next_step_texturePjjjj,480.32,4,0.5,64,64,856,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +12293729a068f060,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +12293729a069c6a0,_Z29kernel_levelines_texture_smemPjjjj,2676.93,3,0.5,64,64,2621500,8,8,1,0,584,18,0,-1,,0,47256,10,783580,386407,409,0,0,0,0,0,12788,0,0,0,270,1214898,26804 +12293729a092a9c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +12293729a0938000,_Z30kernel_estim_next_step_texturePjjjj,475.328,3,0.5,64,64,32767,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +12293729a09acb00,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +12293729a09ba140,_Z29kernel_levelines_texture_smemPjjjj,2666.72,4,0.5,64,64,0,8,8,1,0,584,18,0,-1,,0,47606,13,789385,389081,410,0,0,0,0,0,12815,0,0,0,272,1215491,28047 +12293729a0c45c60,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +12293729a0c53260,_Z30kernel_estim_next_step_texturePjjjj,479.136,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +12293729a0d07640,memcpyAtoH,11.712,30,,,,,,,,,,,0,0,576,0 +12293729a0d10ae0,memcpyDtoH,191.936,733,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_smem_3staticArrays_Context_0.csv b/lniv_smem_3staticArrays_Context_0.csv new file mode 100644 index 0000000..a1b06dd --- /dev/null +++ b/lniv_smem_3staticArrays_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef8902a258 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122935912c8a99a0,memcpyHtoA,177.376,332,,,,,,,,,,,0,0,1048576,0 +122935912c911b40,_Z19kernel_calcul_pathsP4int2j,42.752,12,0.031,1,1,11338272,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122935912c91cca0,memcpyDtoA,4.32,5,,,,,,,,,,,0,0,576,0 +122935912c921a40,_Z23kernel_neutre_img2estimPjjj,440.864,9,0.5,64,64,1074266112,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122935912c9973a0,memcpyAtoD,50.688,6,,,,,,,,,,,0,0,1048576,0 +122935912c9a4400,memcpyDtoA,47.392,3,,,,,,,,,,,0,0,1048576,0 +122935912c9b5440,_Z29kernel_levelines_texture_smemPjjjj,3593.66,9,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,72362,4638,841352,995747,410,0,0,0,0,0,12802,0,0,0,272,1199518,41186 +122935912cd23460,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122935912cd32a20,_Z30kernel_estim_next_step_texturePjjjj,478.112,6,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935912cda8020,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935912cdb5600,_Z29kernel_levelines_texture_smemPjjjj,3592.77,4,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72795,5071,843086,996169,410,0,0,0,0,0,12804,0,0,0,272,1199848,42008 +122935912d1232c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912d130960,_Z30kernel_estim_next_step_texturePjjjj,463.392,4,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935912d1a25e0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122935912d1afc40,_Z29kernel_levelines_texture_smemPjjjj,3594.72,3,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72875,5149,843404,1000954,409,0,0,0,0,0,12800,0,0,0,272,1195816,44504 +122935912d51e060,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912d52b660,_Z30kernel_estim_next_step_texturePjjjj,475.52,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935912d5a0220,memcpyDtoA,47.552,11,,,,,,,,,,,0,0,1048576,0 +122935912d5ad800,_Z29kernel_levelines_texture_smemPjjjj,3590.11,4,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72741,5017,842871,999358,409,0,0,0,0,0,12788,0,0,0,272,1197057,43263 +122935912d91aa40,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912d928060,_Z30kernel_estim_next_step_texturePjjjj,489.248,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122935912d9a01c0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912d9ad840,_Z29kernel_levelines_texture_smemPjjjj,3597.86,4,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72787,5061,843057,1001133,410,0,0,0,0,0,12815,0,0,0,272,1198644,43596 +122935912dd1c8a0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935912dd29f00,_Z30kernel_estim_next_step_texturePjjjj,481.728,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935912dda0320,memcpyDtoA,47.68,3,,,,,,,,,,,0,0,1048576,0 +122935912ddad980,_Z29kernel_levelines_texture_smemPjjjj,3595.36,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,72422,4944,839570,998486,410,0,0,0,0,0,12802,0,0,0,271,1198607,42097 +122935912e11c040,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122935912e129680,_Z30kernel_estim_next_step_texturePjjjj,477.28,3,0.5,64,64,18,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122935912e19e940,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935912e1abf20,_Z29kernel_levelines_texture_smemPjjjj,3600.32,3,0.5,64,64,10083328,8,8,1,0,1064,17,0,-1,,0,72371,4891,839366,996160,410,0,0,0,0,0,12804,0,0,0,271,1199727,42129 +122935912e51b920,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122935912e528f80,_Z30kernel_estim_next_step_texturePjjjj,474.72,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935912e59d840,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935912e5aae60,_Z29kernel_levelines_texture_smemPjjjj,3596.1,3,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72068,4840,836118,990542,409,0,0,0,0,0,12800,0,0,0,270,1197654,42666 +122935912e919820,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122935912e926e80,_Z30kernel_estim_next_step_texturePjjjj,476.032,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935912e99bc60,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935912e9a9280,_Z29kernel_levelines_texture_smemPjjjj,3600.45,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,72013,4785,835903,990163,409,0,0,0,0,0,12788,0,0,0,270,1199082,41238 +122935912ed18d20,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122935912ed26360,_Z30kernel_estim_next_step_texturePjjjj,479.328,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15625,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935912ed9be20,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122935912eda9440,_Z29kernel_levelines_texture_smemPjjjj,3593.73,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,71997,4769,835835,989090,410,0,0,0,0,0,12815,0,0,0,270,1199031,43209 +122935912f117480,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122935912f124a60,_Z30kernel_estim_next_step_texturePjjjj,474.816,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122935912f1993a0,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122935912f1a6a00,_Z29kernel_levelines_texture_smemPjjjj,3588.06,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,71962,4734,835695,988526,410,0,0,0,0,0,12802,0,0,0,270,1197717,42987 +122935912f513440,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122935912f520aa0,_Z30kernel_estim_next_step_texturePjjjj,492.384,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122935912f599860,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122935912f5a6ea0,_Z29kernel_levelines_texture_smemPjjjj,3582.43,4,0.5,64,64,-1386403272,8,8,1,0,1064,17,0,-1,,0,71991,4763,835813,989348,410,0,0,0,0,0,12804,0,0,0,270,1198967,42889 +122935912f912300,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912f91f940,_Z30kernel_estim_next_step_texturePjjjj,476.928,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935912f994aa0,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935912f9a20a0,_Z29kernel_levelines_texture_smemPjjjj,3605.02,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,71919,4691,835524,991687,409,0,0,0,0,0,12800,0,0,0,270,1195048,45272 +122935912fd12d20,memcpyDtoA,47.744,3,,,,,,,,,,,0,0,1048576,0 +122935912fd203c0,_Z30kernel_estim_next_step_texturePjjjj,475.52,4,0.5,64,64,32,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122935912fd94f80,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122935912fda2580,_Z29kernel_levelines_texture_smemPjjjj,3586.85,4,0.5,64,64,32,8,8,1,0,1064,17,0,-1,,0,71947,4719,835634,989254,409,0,0,0,0,0,12788,0,0,0,270,1197194,43126 +122935913010eb00,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122935913011c100,_Z30kernel_estim_next_step_texturePjjjj,477.696,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +1229359130191560,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122935913019eba0,_Z29kernel_levelines_texture_smemPjjjj,3588.93,4,0.5,64,64,65280,8,8,1,0,1064,17,0,-1,,0,72420,4696,841586,995223,410,0,0,0,0,0,12815,0,0,0,272,1199000,43240 +122935913050b940,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +1229359130518f60,_Z30kernel_estim_next_step_texturePjjjj,482.272,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +12293591305cd240,memcpyAtoH,11.584,32,,,,,,,,,,,0,0,576,0 +12293591305d6be0,memcpyDtoH,191.232,734,,,,,,,,,,,0,0,1048576,0 diff --git a/lniv_smem_3staticArrays_Context_0.trc b/lniv_smem_3staticArrays_Context_0.trc new file mode 100644 index 0000000..4589c3e Binary files /dev/null and b/lniv_smem_3staticArrays_Context_0.trc differ diff --git a/lniv_smem_3staticArrays_Context_1.trc b/lniv_smem_3staticArrays_Context_1.trc new file mode 100644 index 0000000..f7a3b2e Binary files /dev/null and b/lniv_smem_3staticArrays_Context_1.trc differ diff --git a/lniv_smem_3staticArrays_Context_2.trc b/lniv_smem_3staticArrays_Context_2.trc new file mode 100644 index 0000000..d33cbfa Binary files /dev/null and b/lniv_smem_3staticArrays_Context_2.trc differ diff --git a/lniv_smem_zc_CFI_linearArray_Context_0.csv b/lniv_smem_zc_CFI_linearArray_Context_0.csv new file mode 100644 index 0000000..387d763 --- /dev/null +++ b/lniv_smem_zc_CFI_linearArray_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef8a3fef18 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122937ea2858cb60,memcpyHtoA,214.656,809,,,,,,,,,,,0,0,1048576,0 +122937ea28726900,_Z19kernel_calcul_pathsP4int2j,42.144,16,0.031,1,1,11690128,1,1,1,0,28,11,0,-1,,0,168,0,1588,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122937ea287317e0,memcpyDtoA,4.128,7,,,,,,,,,,,0,0,576,0 +122937ea28736640,_Z23kernel_neutre_img2estimPjjj,453.632,10,0.5,64,64,4,8,8,1,0,32,5,0,-1,,0,0,0,5496,0,410,0,0,0,0,0,13120,0,0,0,274,2460,820 +122937ea288f0a20,memcpyAtoD,51.296,7,,,,,,,,,,,0,0,1048576,0 +122937ea288fdce0,memcpyDtoA,47.456,4,,,,,,,,,,,0,0,1048576,0 +122937ea2890ef80,_Z29kernel_levelines_texture_smemPjjjj,2313.95,11,0.5,64,64,11454132,8,8,1,0,308,18,0,-1,,0,47434,12,760498,69890,410,0,0,0,0,0,12760,0,0,0,271,1211981,29748 +122937ea28b448e0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122937ea28b53ec0,_Z30kernel_estim_next_step_texturePjjjj,497.632,7,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15636,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea28bce100,memcpyDtoA,47.712,4,,,,,,,,,,,0,0,1048576,0 +122937ea28bdb740,_Z29kernel_levelines_texture_smemPjjjj,2293.38,4,0.5,64,64,337,8,8,1,0,308,18,0,-1,,0,47434,12,760497,69704,410,0,0,0,0,0,12760,0,0,0,271,1216149,27759 +122937ea28e0c000,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937ea28e195c0,_Z30kernel_estim_next_step_texturePjjjj,478.784,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea28e8ee40,memcpyDtoA,1401.6,4,,,,,,,,,,,0,0,1048576,0 +122937ea28fe6f20,_Z29kernel_levelines_texture_smemPjjjj,2285.98,3,0.5,64,64,40,8,8,1,0,308,18,0,-1,,0,47434,12,760495,69267,410,0,0,0,0,0,12760,0,0,0,271,1211639,27223 +122937ea29215b20,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937ea292230e0,_Z30kernel_estim_next_step_texturePjjjj,476.352,4,0.5,64,64,127,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea29298000,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122937ea292a5580,_Z29kernel_levelines_texture_smemPjjjj,2276.77,4,0.5,64,64,8415296,8,8,1,0,308,18,0,-1,,0,47434,12,760494,68848,410,0,0,0,0,0,12760,0,0,0,271,1202995,26669 +122937ea294d1d60,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937ea294df340,_Z30kernel_estim_next_step_texturePjjjj,480.416,4,0.5,64,64,1072693248,8,8,1,0,36,9,0,-1,,0,822,0,15633,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea29555240,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937ea296a25a0,_Z29kernel_levelines_texture_smemPjjjj,2269.89,4,0.5,64,64,1,8,8,1,0,308,18,0,-1,,0,47434,12,760496,68781,410,0,0,0,0,0,12760,0,0,0,271,1203474,26190 +122937ea298cd2a0,memcpyDtoA,47.744,3,,,,,,,,,,,0,0,1048576,0 +122937ea298da900,_Z30kernel_estim_next_step_texturePjjjj,486.432,3,0.5,64,64,16,8,8,1,0,36,9,0,-1,,0,822,0,15635,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea29951f60,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +122937ea2995f5a0,_Z29kernel_levelines_texture_smemPjjjj,2277.66,4,0.5,64,64,0,8,8,1,0,308,18,0,-1,,0,47434,12,760500,68764,410,0,0,0,0,0,12760,0,0,0,271,1205094,24836 +122937ea29b8c100,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937ea29b996e0,_Z30kernel_estim_next_step_texturePjjjj,491.712,4,0.5,64,64,8,8,8,1,0,36,9,0,-1,,0,822,0,15636,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea29c12200,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122937ea29d1e8a0,_Z29kernel_levelines_texture_smemPjjjj,2273.31,4,0.5,64,64,8755104,8,8,1,0,308,18,0,-1,,0,47434,12,760499,68617,410,0,0,0,0,0,12760,0,0,0,271,1202684,27194 +122937ea29f4a300,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122937ea29f578e0,_Z30kernel_estim_next_step_texturePjjjj,473.216,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea29fcbbc0,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937ea29fd91a0,_Z29kernel_levelines_texture_smemPjjjj,2265.38,4,0.5,64,64,0,8,8,1,0,308,18,0,-1,,0,47434,12,760499,68678,410,0,0,0,0,0,12760,0,0,0,271,1205362,24302 +122937ea2a202d00,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122937ea2a210320,_Z30kernel_estim_next_step_texturePjjjj,475.552,3,0.5,64,64,1076494336,8,8,1,0,36,9,0,-1,,0,822,0,15636,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2a284f00,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937ea2a38c880,_Z29kernel_levelines_texture_smemPjjjj,2268.16,3,0.5,64,64,11096400,8,8,1,0,308,18,0,-1,,0,47434,12,760496,68770,410,0,0,0,0,0,12760,0,0,0,271,1205150,24514 +122937ea2a5b6ec0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122937ea2a5c44c0,_Z30kernel_estim_next_step_texturePjjjj,474.496,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2a638ca0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122937ea2a6462e0,_Z29kernel_levelines_texture_smemPjjjj,2277.92,4,0.5,64,64,144,8,8,1,0,308,18,0,-1,,0,47434,12,760500,68681,410,0,0,0,0,0,12760,0,0,0,271,1204010,25654 +122937ea2a872f80,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937ea2a880540,_Z30kernel_estim_next_step_texturePjjjj,472.32,3,0.5,64,64,48,8,8,1,0,36,9,0,-1,,0,822,0,15636,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2a8f4480,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122937ea2aa00920,_Z29kernel_levelines_texture_smemPjjjj,2275.33,4,0.5,64,64,11937498,8,8,1,0,308,18,0,-1,,0,47606,11,763277,68860,410,0,0,0,0,0,12760,0,0,0,272,1203314,25998 +122937ea2ac2cb80,memcpyDtoA,47.776,3,,,,,,,,,,,0,0,1048576,0 +122937ea2ac3a1e0,_Z30kernel_estim_next_step_texturePjjjj,479.136,4,0.5,64,64,6357093,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2acafbc0,memcpyDtoA,47.648,3,,,,,,,,,,,0,0,1048576,0 +122937ea2acbd1a0,_Z29kernel_levelines_texture_smemPjjjj,2268.19,5,0.5,64,64,6488165,8,8,1,0,308,18,0,-1,,0,47434,12,760497,68864,410,0,0,0,0,0,12760,0,0,0,271,1204950,24714 +122937ea2aee7800,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122937ea2aef4da0,_Z30kernel_estim_next_step_texturePjjjj,468.896,3,0.5,64,64,3211309,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2af67fc0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122937ea2b06fe40,_Z29kernel_levelines_texture_smemPjjjj,2280.93,4,0.5,64,64,3670061,8,8,1,0,308,18,0,-1,,0,47434,12,760497,68707,410,0,0,0,0,0,12760,0,0,0,271,1203092,26572 +122937ea2b29d660,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122937ea2b2aad60,_Z30kernel_estim_next_step_texturePjjjj,486.016,4,0.5,64,64,3407923,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2b322220,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122937ea2b32f800,_Z29kernel_levelines_texture_smemPjjjj,2272.99,4,0.5,64,64,3211314,8,8,1,0,308,18,0,-1,,0,47434,12,760500,69017,410,0,0,0,0,0,12760,0,0,0,271,1205428,24236 +122937ea2b55b120,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122937ea2b5686c0,_Z30kernel_estim_next_step_texturePjjjj,490.4,4,0.5,64,64,3145773,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2b5e0ce0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122937ea2b6e5e20,_Z29kernel_levelines_texture_smemPjjjj,2271.9,3,0.5,64,64,11665408,8,8,1,0,308,18,0,-1,,0,47434,12,760503,68833,410,0,0,0,0,0,12760,0,0,0,271,1203294,26370 +122937ea2b911340,memcpyDtoA,47.712,4,,,,,,,,,,,0,0,1048576,0 +122937ea2b91e980,_Z30kernel_estim_next_step_texturePjjjj,479.776,3,0.5,64,64,10922,8,8,1,0,36,9,0,-1,,0,822,0,15634,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122937ea2ba8b1a0,memcpyAtoH,11.936,782,,,,,,,,,,,0,0,576,0 +122937ea2beeee40,memcpyDtoH,192.032,1729,,,,,,,,,,,0,0,1048576,0 diff --git a/main.cu b/main.cu index d9db4f8..0bb7c3d 100644 --- a/main.cu +++ b/main.cu @@ -13,127 +13,76 @@ // lib spec #include "defines.h" #include "levelines_common.h" - #include "levelines_kernels.cu" -__global__ void kernel_debil(unsigned int * ptr1, unsigned int * ptr2, unsigned int L, int val){ - - unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; - unsigned int j = blockIdx.y*blockDim.y + threadIdx.y; - unsigned int pos = i*L +j ; +int main(int argc, char **argv){ - ptr2[pos] = val - ptr1[pos] ; + // use device with highest Gflops/s + cudaSetDevice( 0 ); -} - -int main(int argc, char **argv){ + unsigned int timer ; + cutilCheckError( cutCreateTimer(&timer) ); + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + + //alloc bidon pour anticiper l'initialisation du GPU + short * d_bidon ; + cutilSafeCall( cudaMalloc( (void**) &d_bidon, sizeof(short))) ; //float coef_regul = atof( argv[1] ) ; - unsigned int timer ; - cutilCheckError( cutCreateTimer(&timer) ); - cutilCheckError( cutResetTimer(timer) ); /***************************** * CHARGEMENT IMAGE *****************************/ char* image_path = argv[argc-1]; - char* image_out = "./image_out.pgm" ; + unsigned int r = atoi(argv[1]) ; + bool seq_out = atoi(argv[2]) ; + unsigned int iter , nb_iter = atoi(argv[3]) ; + unsigned int poids = 15 ; + char * image_out_base = "./image_out" ; + char * pgm_ext = ".pgm" ; + char image_out[80] ; unsigned int * h_data = NULL ; unsigned int * h_data_out = NULL ; - unsigned int H, L, size; - - cutilCheckError( cutStartTimer(timer) ); - cutilCheckError( cutLoadPGMi(image_path, &h_data, &L, &H)); - cutilCheckError( cutStopTimer(timer) ); + unsigned int H, L, size, psize ; - size = H * L * sizeof( unsigned int ); - printf("Loaded %d x %d = %d pixels from '%s' en %f ms,\n", L, H, size, image_path, cutGetTimerValue(timer)); - - - //essai alloc mapped - /* - cutilCheckError( cutResetTimer(timer) ); - cutilCheckError( cutStartTimer(timer) ); - unsigned int * h_ptr1, * d_ptr1 ; - unsigned int * h_ptr2, * d_ptr2 ; - int h = ; - int l = h ; - int mem = h*l*sizeof(unsigned int) ; - cutilSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost)); - cutilCheckError( cutStopTimer(timer) ); - printf("Temps set flag Mapped : %f ms\n", cutGetTimerValue(timer)) ; - - cutilCheckError( cutStartTimer(timer) ); - cutilSafeCall(cudaHostAlloc((void **)&h_ptr1, mem, cudaHostAllocMapped)); - cutilSafeCall(cudaHostAlloc((void **)&h_ptr2, mem, cudaHostAllocMapped)); + // chargt image + cutilCheckError( cutLoadPGMi(image_path, &h_data, &L, &H)); cutilCheckError( cutStopTimer(timer) ); - printf("Temps cumul alloc Mapped : %f ms\n", cutGetTimerValue(timer)) ; - for (int i = 0; i>>(d_ptr1, d_ptr2, l, 255) ; + printf("Longueur des chemins = %d pixels\n", r); + printf("Init GPU + Image %s %d x %d = %d pixels en %f ms,\n", image_path, L, H, size, cutGetTimerValue(timer)); - cutilCheckError( cutStopTimer(timer) ); - printf("Temps total Mapped : %f ms\n", cutGetTimerValue(timer)) ; - - char * image_1 = "./image_1.pgm" ; - char * image_2 = "./image_2.pgm" ; - - cutilCheckError( cutSavePGMi(image_1, h_ptr1, l, h) ) ; - cutilCheckError( cutSavePGMi(image_2, h_ptr2, l, h) ) ; - */ /***************************** * FIN CHARGEMENT IMAGE *****************************/ - - - - // use device with highest Gflops/s - cudaSetDevice( cutGetMaxGflopsDeviceId() ); - - - /* - cutilSafeCall( cudaMallocArray(&a_Src, &floatTex, imageW, imageH) ); - cutilSafeCall( cudaMalloc((void **)&d_Output, imageW * imageH * sizeof(float)) ); - cutilSafeCall( cudaThreadSynchronize() ); - cutilCheckError( cutResetTimer(hTimer) ); - cutilCheckError( cutStartTimer(hTimer) ); - - cutilSafeCall( cudaThreadSynchronize() ); - cutilCheckError( cutStopTimer(hTimer) ); - gpuTime = cutGetTimerValue(hTimer) / (float)iterations; - */ cutilCheckError( cutResetTimer(timer) ); cutilCheckError( cutStartTimer(timer) ); // allocation mem GPU unsigned int * d_directions =NULL ; unsigned int * d_lniv, * d_estim = NULL ; + int2 * d_paths ; cutilSafeCall( cudaMalloc( (void**) &d_directions, size)) ; cutilSafeCall( cudaMalloc( (void**) &d_lniv, size ) ); cutilSafeCall( cudaMalloc( (void**) &d_estim, size ) ); + cutilSafeCall( cudaMalloc( (void**) &d_paths, psize ) ); // allocate array and copy image data cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindUnsigned); - cudaArray * array_img_in, *array_img_estim, *array_img_lniv; + cudaChannelFormatDesc channelDescP = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindSigned); + + cudaArray * array_img_in, *array_img_estim, *array_img_lniv, *array_paths; cutilSafeCall( cudaMallocArray( &array_img_in, &channelDesc, L, H )); cutilSafeCall( cudaMemcpyToArray( array_img_in, 0, 0, h_data, size, cudaMemcpyHostToDevice)) ; cutilSafeCall( cudaBindTextureToArray( tex_img_in, array_img_in, channelDesc)); - cutilCheckError( cutStopTimer(timer) ); cutilSafeCall( cudaMallocArray( &array_img_estim, &channelDesc, L, H )); cutilSafeCall( cudaBindTextureToArray( tex_img_estim, array_img_estim, channelDesc)); @@ -141,28 +90,75 @@ int main(int argc, char **argv){ cutilSafeCall( cudaMallocArray( &array_img_lniv, &channelDesc, L, H )); cutilSafeCall( cudaBindTextureToArray( tex_img_lniv, array_img_lniv, channelDesc)); + cutilSafeCall( cudaMallocArray( &array_paths, &channelDescP, (r-1), PSIZE_I )); + cutilSafeCall( cudaBindTextureToArray( tex_paths, array_paths, channelDescP)); + + cutilCheckError( cutStopTimer(timer) ); printf("Temps alloc + transferts en Textures : %f ms\n", cutGetTimerValue(timer)) ; + /***************************** - * APPELS KERNELS et chronos + * GENERATION des CHEMINS *****************************/ cutilCheckError( cutResetTimer(timer) ); cutilCheckError( cutStartTimer(timer) ); - unsigned int iter , nb_iter = 15 ; - unsigned int poids = 15 ; - dim3 dimBlock(8,8,1) ; - dim3 dimGrid( H / dimBlock.x, L / dimBlock.y, 1 ) ; - unsigned int smem_size = dimBlock.x * dimBlock.y * sizeof(unsigned int) ; - // init image estimee avec image_in + dim3 dimBlock(1,1,1) ; + dim3 dimGrid(1,1,1) ; + // calcul des chemins + kernel_calcul_paths<<< dimGrid, dimBlock, 0 >>>(d_paths, r); + + // copie du tableau en texture + cutilSafeCall( cudaMemcpyToArray( array_paths, 0, 0, d_paths, psize, cudaMemcpyDeviceToDevice)) ; + + cutilCheckError( cutStopTimer(timer) ); + printf("Temps generation chemin + transfert en texture : %f ms\n", cutGetTimerValue(timer)) ; + + /***************************** + * APPELS KERNELS et chronos + *****************************/ + dimBlock = dim3(16,16,1) ; + dimGrid = dim3( H / dimBlock.x,L / dimBlock.y, 1 ) ; + + // pour enregistrement image lniv GPU + free(h_data_out) ; + h_data_out = new unsigned int[H*L] ; + + //init image estimee avec image_in + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); kernel_init_estim_from_img_in<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, 7); - printf("Grille : %d x %d de Blocs : %d x %d - Shared mem : %d octets\n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y, smem_size) ; + // kernel de copie + /* + kernel_neutre_img2estim<<< dimGrid, dimBlock, 0>>>(d_estim, L, H); + */ + cudaThreadSynchronize() ; + cutilCheckError( cutStopTimer(timer) ); + printf("Temps kernel init : %f ms\n", cutGetTimerValue(timer)) ; + // a remplacer par + /* + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + cutilSafeCall( cudaMemcpyFromArray( d_estim, array_img_in, 0, 0, size, cudaMemcpyDeviceToDevice)) ; + cutilCheckError( cutStopTimer(timer) ); + printf("Temps memcpyFromArray : %f ms\n", cutGetTimerValue(timer)) ; + */ + printf("Grille : %d x %d de Blocs : %d x %d \n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y) ; + + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); for ( iter =0 ; iter < nb_iter ; iter++ ) { cutilSafeCall( cudaMemcpyToArray( array_img_estim, 0, 0, d_estim, size, cudaMemcpyDeviceToDevice)) ; - kernel_levelines_texture<<< dimGrid, dimBlock, 0 >>>( d_lniv, L, H ); + kernel_levelines_texture_smem<<< dimGrid, dimBlock, 0 >>>( d_lniv, L, H, r ); cutilSafeCall( cudaMemcpyToArray( array_img_lniv, 0, 0, d_lniv, size, cudaMemcpyDeviceToDevice)) ; + if (seq_out){ + sprintf(image_out, "%s%d%s", image_out_base, iter+1, pgm_ext) ; + printf("chaine : %s\n", image_out); + cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) ); + cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ; + } kernel_estim_next_step_texture<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, poids) ; } @@ -171,26 +167,95 @@ int main(int argc, char **argv){ cutilCheckError( cutStopTimer(timer) ); printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)/(float)nb_iter) ; printf("Total pour %d kernels : %f ms\n", nb_iter, cutGetTimerValue(timer)) ; - + /************************** * VERIFS **************************/ - //trace des lniv sur grille de 'pas x pas' - //kernel_trace_levelines<<< dimGrid, dimBlock, 0 >>>(d_data, d_directions, d_data2, L, H, 16, 255) ; - //cudaThreadSynchronize(); - - // enregistrement image lniv GPU - h_data_out = new unsigned int[H*L] ; - if ( h_data_out != NULL) - cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) ); - else - printf("Echec allocation mem CPU\n"); - cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ; + /************************************************** + * recuperation matrice des chemins pour affichage + **************************************************/ + + int2 * h_paths = new int2[(r-1)*PSIZE_I] ; + cutilSafeCall( cudaMemcpyFromArray(h_paths , array_paths, 0, 0, psize, cudaMemcpyDeviceToHost) ); + /* + //verif Di + printf("matrice Di"); + for(int idpath=0; idpath< PSIZE_I; idpath++){b + + printf("\n"); + for(int idpix=0; idpix< r-1; idpix++){ + printf(" % d ", h_paths[idpath*(r-1) + idpix].x ); + } + printf("\t// %d°", idpath*15) ; + } + //verif Dj + printf("\nmatrice Dj"); + for(int idpath=0; idpath< PSIZE_I; idpath++){ + printf("\n"); + for(int idpix=0; idpix< r-1; idpix++){ + printf(" % d ", h_paths[idpath*(r-1) + idpix].y); + } + printf("\t// %d°", idpath*15) ; + } + printf("\n"); + */ + /*************************************************** + * fin verif visuelle matrices des chemins + ***************************************************/ - // calcul lniv CPU + /*************************************************** + * execution sequentielle pour comparaison + * la comparaison n'est pertinente que + * si d_lniv contient les lniv de l'image se départ + **************************************************/ + /* + // calcul sequentiel + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + h_data_out = new unsigned int[H*L] ; + int * dout = new int[H*L] ; + for ( iter =0 ; iter < nb_iter ; iter++ ){ + for (int i=r; i<= H-r; i++){ + for (int j=r; j<= L-r; j++){ + h_data_out[i*L+j] = lniv4_value(h_data, h_paths, i, j, H, L, &dout[i*L+j], r) ; + } + } + } + cutilCheckError( cutStopTimer(timer) ); + printf("Execution sequentielle CPU : %f ms\n", cutGetTimerValue(timer)) ; + // comparaison + unsigned int * h_lniv = new unsigned int[H*L] ; + int pos, cpt_err=0, cpt_pix=0 ; + cutilSafeCall( cudaMemcpy(h_lniv , d_lniv, size, cudaMemcpyDeviceToHost) ); + for ( iter =0 ; iter < nb_iter ; iter++ ){ + for (int i=r; i<= H-r; i++){ + for (int j=r; j<= L-r; j++){ + pos = i*L + j ; + if ( h_data_out[ pos ] != h_lniv[ pos ] ) { + cpt_err++ ; + printf(" pixel ( %d , %d ) -> GPU= %d CPU= %d \n", i, j, h_lniv[pos], h_data_out[pos]); + } + cpt_pix++ ; + } + } + } + printf("TAUX ERREUR GPU/CPU : %d / %d \n", cpt_err, cpt_pix ); + */ + //trace des lniv sur grille de 'pas x pas' + //kernel_trace_levelines<<< dimGrid, dimBlock, 0 >>>(d_data, d_directions, d_data2, L, H, 16, 255, r) ; + + if (!seq_out){ + if ( h_data_out != NULL) + cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) ); + else + printf("Echec allocation mem CPU\n"); + sprintf(image_out, "%s%d%s", image_out_base, iter+1, pgm_ext) ; + cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ; + } + // TODO verifier pourquoi les deux lignes suivantes produisent une erreur //cutilExit(argc, argv); //cudaThreadExit(); diff --git a/profile.cvp b/profile.cvp new file mode 100644 index 0000000..f5adf09 --- /dev/null +++ b/profile.cvp @@ -0,0 +1,22 @@ + + + + Session1 + "/home/perrot/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/lniv" + 4 0 15 /home/perrot/Images/cochons/cochon_512b.pgm + /home/perrot/NVIDIA_GPU_Computing_SDK/C/src/lniv_gpu + 21 Jun 2011 10:18:11 + false + + + branch,divergent_branch,instructions,warp_serialize + cta_launched,local_load,local_store,gld_32b + gld_64b,gld_128b,gst_32b,gst_64b + gst_128b + gld_request,gst_request + tex_cache_hit,tex_cache_miss + + + + + diff --git a/profile_Session1_Context_0.csv b/profile_Session1_Context_0.csv new file mode 100644 index 0000000..8f4a1d0 --- /dev/null +++ b/profile_Session1_Context_0.csv @@ -0,0 +1,72 @@ +# CUDA_PROFILE_LOG_VERSION 2.0 +# CUDA_PROFILE_CSV 1 +# TIMESTAMPFACTOR fffff6ef886fccd8 +# CUDA_DEVICE 0 Tesla C1060 +gpustarttimestamp,method,gputime,cputime,occupancy,gridSizeX,gridSizeY,gridSizeZ,blockSizeX,blockSizeY,blockSizeZ,dynSmemPerBlock,staSmemPerBlock,registerPerThread,streamID,localblocksize,memTransferSize,memtransferhostmemtype,branch,divergent_branch,instructions,warp_serialize,cta_launched,local_load,local_store,gld_32b,gld_64b,gld_128b,gst_32b,gst_64b,gst_128b,gld_request,gst_request,tex_cache_hit,tex_cache_miss +122933fc4123b400,memcpyHtoA,177.376,329,,,,,,,,,,,0,0,1048576,0 +122933fc412a2780,_Z19kernel_calcul_pathsP4int2j,43.296,12,0.031,1,1,-1386402296,1,1,1,0,28,11,0,-1,,0,168,0,1587,0,1,0,0,54,0,0,72,0,0,54,72,0,0 +122933fc412adae0,memcpyDtoA,4.256,5,,,,,,,,,,,0,0,576,0 +122933fc412b2800,_Z23kernel_neutre_img2estimPjjj,446.912,9,0.5,64,64,1074266112,8,8,1,0,32,5,0,-1,,0,0,0,5448,0,409,0,0,0,0,0,13088,0,0,0,272,2454,818 +122933fc4132a1e0,memcpyAtoD,50.528,6,,,,,,,,,,,0,0,1048576,0 +122933fc41337180,memcpyDtoA,47.424,4,,,,,,,,,,,0,0,1048576,0 +122933fc41348200,_Z29kernel_levelines_texture_smemPjjjj,3595.01,10,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,72362,4638,841352,997341,410,0,0,0,0,0,12802,0,0,0,272,1196850,43854 +122933fc416b6760,memcpyDtoA,47.744,3,,,,,,,,,,,0,0,1048576,0 +122933fc416c5d80,_Z30kernel_estim_next_step_texturePjjjj,477.824,6,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122933fc4173b240,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122933fc41748840,_Z29kernel_levelines_texture_smemPjjjj,3596.54,4,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72795,5071,843084,998235,410,0,0,0,0,0,12804,0,0,0,272,1199869,41987 +122933fc41ab7380,memcpyDtoA,47.712,3,,,,,,,,,,,0,0,1048576,0 +122933fc41ac4a00,_Z30kernel_estim_next_step_texturePjjjj,480.928,4,0.5,64,64,1075052544,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122933fc41b3aae0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122933fc41b480e0,_Z29kernel_levelines_texture_smemPjjjj,3589.54,4,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72875,5149,843406,999080,409,0,0,0,0,0,12800,0,0,0,272,1195950,44370 +122933fc41eb50e0,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122933fc41ec26e0,_Z30kernel_estim_next_step_texturePjjjj,477.792,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15513,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122933fc41f37bc0,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122933fc41f45200,_Z29kernel_levelines_texture_smemPjjjj,3592.61,3,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72741,5017,842869,995949,409,0,0,0,0,0,12788,0,0,0,272,1196521,43799 +122933fc422b2e00,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122933fc422c03e0,_Z30kernel_estim_next_step_texturePjjjj,461.92,4,0.5,64,64,1074266112,8,8,1,0,36,9,0,-1,,0,816,0,15512,0,410,0,0,0,0,0,13120,0,0,0,272,4920,1640 +122933fc42331a80,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122933fc4233f100,_Z29kernel_levelines_texture_smemPjjjj,3582.94,3,0.5,64,64,1,8,8,1,0,1064,17,0,-1,,0,72787,5061,843054,1000048,410,0,0,0,0,0,12815,0,0,0,272,1197435,44805 +122933fc426aa720,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122933fc426b7da0,_Z30kernel_estim_next_step_texturePjjjj,479.776,4,0.5,64,64,3080302,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122933fc4272da40,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122933fc4273b040,_Z29kernel_levelines_texture_smemPjjjj,3586.02,4,0.5,64,64,3080312,8,8,1,0,1064,17,0,-1,,0,72422,4944,839570,994815,410,0,0,0,0,0,12802,0,0,0,271,1198152,42552 +122933fc42aa7280,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122933fc42ab4880,_Z30kernel_estim_next_step_texturePjjjj,480.544,3,0.5,64,64,7536737,8,8,1,0,36,9,0,-1,,0,816,0,15510,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122933fc42b2a7e0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122933fc42b37de0,_Z29kernel_levelines_texture_smemPjjjj,3599.3,4,0.5,64,64,7733353,8,8,1,0,1064,17,0,-1,,0,72371,4891,839365,996460,410,0,0,0,0,0,12804,0,0,0,271,1199816,42040 +122933fc42ea73e0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122933fc42eb49e0,_Z30kernel_estim_next_step_texturePjjjj,474.464,4,0.5,64,64,3211296,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc42f291a0,memcpyDtoA,47.488,3,,,,,,,,,,,0,0,1048576,0 +122933fc42f36760,_Z29kernel_levelines_texture_smemPjjjj,3595.94,4,0.5,64,64,7143535,8,8,1,0,1064,17,0,-1,,0,72068,4840,836122,989513,409,0,0,0,0,0,12800,0,0,0,270,1197365,42955 +122933fc432a5060,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122933fc432b2700,_Z30kernel_estim_next_step_texturePjjjj,472.832,4,0.5,64,64,7471218,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc43326880,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122933fc43333e80,_Z29kernel_levelines_texture_smemPjjjj,3590.69,4,0.5,64,64,6357101,8,8,1,0,1064,17,0,-1,,0,72013,4785,835899,989204,409,0,0,0,0,0,12788,0,0,0,270,1198226,42094 +122933fc436a1300,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122933fc436ae940,_Z30kernel_estim_next_step_texturePjjjj,496.352,4,0.5,64,64,7274595,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc43728680,memcpyDtoA,47.552,3,,,,,,,,,,,0,0,1048576,0 +122933fc43735c80,_Z29kernel_levelines_texture_smemPjjjj,3592.03,4,0.5,64,64,3080307,8,8,1,0,1064,17,0,-1,,0,71997,4769,835835,989426,410,0,0,0,0,0,12815,0,0,0,270,1196748,45492 +122933fc43aa3640,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122933fc43ab0cc0,_Z30kernel_estim_next_step_texturePjjjj,484.352,4,0.5,64,64,7209071,8,8,1,0,36,9,0,-1,,0,822,0,15624,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122933fc43b27b20,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122933fc43b35180,_Z29kernel_levelines_texture_smemPjjjj,3600.42,4,0.5,64,64,3014754,8,8,1,0,1064,17,0,-1,,0,71962,4734,835697,990688,410,0,0,0,0,0,12802,0,0,0,270,1198105,42599 +122933fc43ea4be0,memcpyDtoA,47.584,4,,,,,,,,,,,0,0,1048576,0 +122933fc43eb21e0,_Z30kernel_estim_next_step_texturePjjjj,477.504,4,0.5,64,64,7471136,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,409,0,0,0,0,0,13088,0,0,0,274,4908,1636 +122933fc43f27580,memcpyDtoA,47.584,3,,,,,,,,,,,0,0,1048576,0 +122933fc43f34ba0,_Z29kernel_levelines_texture_smemPjjjj,3592.61,4,0.5,64,64,2097207,8,8,1,0,1064,17,0,-1,,0,71991,4763,835810,989866,410,0,0,0,0,0,12804,0,0,0,270,1198680,43176 +122933fc442a27a0,memcpyDtoA,47.616,3,,,,,,,,,,,0,0,1048576,0 +122933fc442afe20,_Z30kernel_estim_next_step_texturePjjjj,469.472,4,0.5,64,64,51,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc44323240,memcpyDtoA,47.68,4,,,,,,,,,,,0,0,1048576,0 +122933fc44330880,_Z29kernel_levelines_texture_smemPjjjj,3599.94,3,0.5,64,64,-767360128,8,8,1,0,1064,17,0,-1,,0,71919,4691,835524,989067,409,0,0,0,0,0,12800,0,0,0,270,1193858,46462 +122933fc446a0120,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122933fc446ad740,_Z30kernel_estim_next_step_texturePjjjj,473.952,3,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc44721d00,memcpyDtoA,47.648,4,,,,,,,,,,,0,0,1048576,0 +122933fc4472f340,_Z29kernel_levelines_texture_smemPjjjj,3589.06,3,0.5,64,64,-1386403096,8,8,1,0,1064,17,0,-1,,0,71947,4719,835636,987891,409,0,0,0,0,0,12788,0,0,0,270,1198457,41863 +122933fc44a9c160,memcpyDtoA,47.552,4,,,,,,,,,,,0,0,1048576,0 +122933fc44aa9760,_Z30kernel_estim_next_step_texturePjjjj,476.576,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,822,0,15626,0,410,0,0,0,0,0,13120,0,0,0,274,4920,1640 +122933fc44b1e760,memcpyDtoA,47.616,4,,,,,,,,,,,0,0,1048576,0 +122933fc44b2bdc0,_Z29kernel_levelines_texture_smemPjjjj,3590.59,4,0.5,64,64,0,8,8,1,0,1064,17,0,-1,,0,72420,4696,841587,993210,410,0,0,0,0,0,12815,0,0,0,272,1200675,41565 +122933fc44e991e0,memcpyDtoA,47.776,3,,,,,,,,,,,0,0,1048576,0 +122933fc44ea68a0,_Z30kernel_estim_next_step_texturePjjjj,477.408,4,0.5,64,64,0,8,8,1,0,36,9,0,-1,,0,816,0,15511,0,409,0,0,0,0,0,13088,0,0,0,272,4908,1636 +122933fc44f57380,memcpyAtoH,11.552,30,,,,,,,,,,,0,0,576,0 +122933fc44f605e0,memcpyDtoH,191.136,740,,,,,,,,,,,0,0,1048576,0 diff --git a/profile_Session1_Context_0.trc b/profile_Session1_Context_0.trc new file mode 100644 index 0000000..1fa3e00 Binary files /dev/null and b/profile_Session1_Context_0.trc differ diff --git a/profilerapitrace_0.trc b/profilerapitrace_0.trc new file mode 100644 index 0000000..0ada1fc Binary files /dev/null and b/profilerapitrace_0.trc differ