/**edgeDetection.cu**预发布许可证-仅用于工程反馈和测试目的*。不出售。**模型“edgeDetection”的代码生成**模型版本:1.8×Simulink编码器版本:9.4(R2020B)19-MAY-2020年C++源代码:6月16日星期六11:46:16 2020 **目标选择:GRT。TLC*注:GRT包括额外的基础设施和用于原型的仪器*嵌入式金宝app硬件选择:英特尔-X8664(WINDOWS64)*代码生成目标:未指定*验证结果:未运行*/#包括“edgeDetection.h”#包括“edgeDetection_private.h”#包括“math_constants.h”/*局部函数的前向声明*/静态(全局)(启动)边界(512,1)无效edgeDetection(eML)blk(内核)内核(1(const real32(T RGB[230400],real)T gray[76800])/*局部函数的前向声明*/static\uuuuu global\uuuuuu launch\uu bounds\uuuuu(512,1)void edgeDe\u eML\u blk\u kernel\u kernel1\l(real32\u T expanded[77924]);静态全局启动边界(512,1)无效边缘内核2(常量实数灰度图像[76800],实数32扩展[77924]);静态全局启动边界(1024,1)无效边缘内核3(常量int8\T b[9],常量real32\T扩展[77924],real32\T H[76800]);静态、全局、启动、边界(512,1)无效边缘、eML、blk、kernel4(real32扩展[77924]);静态全局启动边界(512,1)无效边缘内核5(常量实数灰度图像[76800],实数32扩展[77924]);静态全局启动边界(1024,1)无效边缘内核(const int8_T b_b[9],const real32_T expanded[77924],real32_T V[76800]);静态全局启动边界(512,1)无效边数据块内核7(常数实阈值,常数实32μT V[76800],uint8μT边数据[76800],实32μT H[76800]);静态全局启动边界(512,1)无效边(eML)blk内核1(const real32_T RGB[230400],real_T gray[76800]){uint64_T threadId;int32_T gray_tmp;threadId=mGetGlobalThreadIndex();gray_tmp=static_cast(threadId%240ULL);threadId=(threadId-static_cast(static_cast(threadId)<320)和(static_cast(gray_tmp<240)){gray_tmp+=240*static_cast(threadId);gray[gray_tmp]=(static_cast(RGB[gray_tmp+76800])*0.587+static_cast(RGB[gray_tmp]=0.2989)+static_cast(RGB[gray]/*MATLAB函数的函数:'/RGB to Gray'*/void edgeDetectionModelClass::edgeDetection_eML_blk_内核(const real32_T RGB[230400],real_T Gray[76800]){real_T(*gpu_Gray)[76800];real32_T(*gpu_RGB)[230400];cudamaloc(&gpu_Gray,614400ULL);cudamaloc(&gpu__RGB,92160ull);cudamcpy,921600ULL,cudamemcpyhostodice);edgeDete_eML_blk_kernel_kernel1<>(*gpu_RGB,*gpu灰);cudaMemcpy(&gray[0],gpu_灰,614400ULL,cudamemcydevicetohost);cudaFree(*gpu RGB);cudaFree(*gpu灰)/*原子系统的输出和更新:'/RGB to Gray'*/void edgeDetectionModelClass::edgeDetection_RGBtoGray(const real32_T rtu_RGB[230400],B_RGBtoGray_edgeDetection_T*localB){edgeDetection_eML_blk_内核(rtu RGB,localB->dv);std::memcpy(&localB->Gray[0],&localB[0],76800U*sizeof(real_T))}静态的全局的线程ID;线程ID=mwGetGlobalThreadIndex();如果(静态的线程ID(threadId%240ULL);threadId=(threadId-static_cast(n))/240ULL;if((static_cast(threadId)<320))&(static_cast(threadId)(threadId)+1))+1]=static_-cast(grayImage[240*static_-cast(threadId)+n]}static_-global_-launch_-bounds_(1024,1)void-edgedetete_-eML_-blk_-k_-kernel_-3(const-int8_-T-T-b[9],const-real32_-T-T-T-T-expanded[77924],real32_-T-T-T-H[77924],real32_-800-u-shared-real6];int32_T baseC;int32_T baseR;int32_T ocol;int32_T orow;int32_T scol;int32_T srow;int32_T stridcol;int32_T strickrow;int32_T x idx;int32_T y idx;real32_T cv;ocol=mGetGlobalThreadIndexinDimension();orow=mGetGlobalThreadIndexinDimension();baseR=orow;srow=static_cast(blockDim.x);scol=static_cast(threadIdx.y);stridcol=static_cast(blockDim.y);for(y_idx=srow;y_idx<=33;y_idx+=stridrow){baseC=ocol;for(x_idx=scol;x_idx<=33;x_idx+=stripcol){if((static_cast((static_cast(baseR >= 0)) && ( static_cast(baseR < 242)))) && (static_cast(( static_cast(baseC >= 0)) && (static_cast(baseC < 322))))) { expanded_shared[y_idx + 34 * x_idx] = expanded[242 * baseC + baseR]; } else { expanded_shared[y_idx + 34 * x_idx] = 0.0F; } baseC += strideCol; } baseR += strideRow; } __syncthreads(); if ((static_cast(ocol < 320)) && (static_cast(orow < 240))) { cv = 0.0F; for (baseR = 0; baseR < 3; baseR++) { strideRow = (baseR + ocol) * 242 + orow; strideCol = (2 - baseR) * 3; cv += expanded_shared[((srow + strideRow % 242) - orow) + 34 * ((scol + strideRow / 242) - ocol)] * static_cast(b[strideCol + 2]); cv += expanded_shared[((srow + (strideRow + 1) % 242) - orow) + 34 * ((scol + (strideRow + 1) / 242) - ocol)] * static_cast (b[strideCol + 1]); cv += expanded_shared[((srow + (strideRow + 2) % 242) - orow) + 34 * ((scol + (strideRow + 2) / 242) - ocol)] * static_cast (b[strideCol]); } H[orow + 240 * ocol] = cv; } } static __global__ __launch_bounds__(512, 1) void edgeDete_eML_blk_kernel_kernel4 (real32_T expanded[77924]) { uint64_T threadId; threadId = mwGetGlobalThreadIndex(); if (static_cast(threadId) < 77924) { expanded[static_cast(threadId)] = 0.0F; } } static __global__ __launch_bounds__(512, 1) void edgeDete_eML_blk_kernel_kernel5 (const real_T grayImage[76800], real32_T expanded[77924]) { uint64_T threadId; int32_T n; threadId = mwGetGlobalThreadIndex(); n = static_cast(threadId % 240ULL); threadId = (threadId - static_cast(n)) / 240ULL; if ((static_cast(static_cast(threadId) < 320)) && ( static_cast(n < 240))) { expanded[(n + 242 * (static_cast(threadId) + 1)) + 1] = static_cast(grayImage[240 * static_cast(threadId) + n]); } } static __global__ __launch_bounds__(1024, 1) void edgeDete_eML_blk_kernel_kernel6(const int8_T b_b[9], const real32_T expanded [77924], real32_T V[76800]) { __shared__ real32_T expanded_shared[1156]; int32_T baseC; int32_T baseR; int32_T ocol; int32_T orow; int32_T scol; int32_T srow; int32_T strideCol; int32_T strideRow; int32_T x_idx; int32_T y_idx; real32_T cv; ocol = mwGetGlobalThreadIndexInYDimension(); orow = mwGetGlobalThreadIndexInXDimension(); baseR = orow; srow = static_cast(threadIdx.x); strideRow = static_cast(blockDim.x); scol = static_cast(threadIdx.y); strideCol = static_cast(blockDim.y); for (y_idx = srow; y_idx <= 33; y_idx += strideRow) { baseC = ocol; for (x_idx = scol; x_idx <= 33; x_idx += strideCol) { if ((static_cast((static_cast(baseR >= 0)) && ( static_cast(baseR < 242)))) && (static_cast(( static_cast(baseC >= 0)) && (static_cast(baseC < 322))))) { expanded_shared[y_idx + 34 * x_idx] = expanded[242 * baseC + baseR]; } else { expanded_shared[y_idx + 34 * x_idx] = 0.0F; } baseC += strideCol; } baseR += strideRow; } __syncthreads(); if ((static_cast(ocol < 320)) && (static_cast(orow < 240))) { cv = 0.0F; for (baseR = 0; baseR < 3; baseR++) { cv += expanded_shared[((srow + orow) - orow) + 34 * (((scol + baseR) + ocol) - ocol)] * static_cast(b_b[(2 - baseR) * 3 + 2]); cv += expanded_shared[(((srow + orow) - orow) + 34 * (((scol + baseR) + ocol) - ocol)) + 1] * static_cast(b_b[(2 - baseR) * 3 + 1]); cv += expanded_shared[(((srow + orow) - orow) + 34 * (((scol + baseR) + ocol) - ocol)) + 2] * static_cast(b_b[(2 - baseR) * 3]); } V[orow + 240 * ocol] = cv; } } static __global__ __launch_bounds__(512, 1) void edgeDete_eML_blk_kernel_kernel7 (const real_T threshold, const real32_T V[76800], uint8_T edgeImage[76800], real32_T H[76800]) { uint64_T threadId; threadId = mwGetGlobalThreadIndex(); if (static_cast(threadId) < 76800) { H[static_cast(threadId)] = H[static_cast(threadId)] * H[ static_cast(threadId)] + V[static_cast(threadId)] * V[ static_cast(threadId)]; H[static_cast(threadId)] = sqrtf(H[static_cast(threadId)]); edgeImage[static_cast(threadId)] = static_cast( static_cast(static_cast(static_cast(H[ static_cast(threadId)]) > threshold)) * 255U); } } /* Function for MATLAB Function: '/Sobel Edge' */ void edgeDetectionModelClass::edgeDetection_eML_blk_kernel_b(const real_T grayImage[76800], real_T threshold, uint8_T edgeImage[76800]) { static const int8_T b[9] = { 1, 0, -1, 2, 0, -2, 1, 0, -1 }; static const int8_T b_b[9] = { 1, 2, 1, 0, 0, 0, -1, -2, -1 }; real_T (*gpu_grayImage)[76800]; real32_T (*gpu_expanded)[77924]; real32_T (*gpu_H)[76800]; real32_T (*gpu_V)[76800]; int8_T (*gpu_b)[9]; int8_T (*gpu_b_b)[9]; uint8_T (*gpu_edgeImage)[76800]; cudaMalloc(&gpu_edgeImage, 76800ULL); cudaMalloc(&gpu_V, 307200ULL); cudaMalloc(&gpu_H, 307200ULL); cudaMalloc(&gpu_b_b, 9ULL); cudaMalloc(&gpu_b, 9ULL); cudaMalloc(&gpu_grayImage, 614400ULL); cudaMalloc(&gpu_expanded, 311696ULL); edgeDe_eML_blk_kernel_kernel1_l<<>> (*gpu_expanded); cudaMemcpy(gpu_grayImage, (void *)&grayImage[0], 614400ULL, cudaMemcpyHostToDevice); edgeDete_eML_blk_kernel_kernel2<<>> (*gpu_grayImage, *gpu_expanded); cudaMemcpy(gpu_b, (void *)&b[0], 9ULL, cudaMemcpyHostToDevice); edgeDete_eML_blk_kernel_kernel3<<>> (*gpu_b, *gpu_expanded, *gpu_H); edgeDete_eML_blk_kernel_kernel4<<>> (*gpu_expanded); edgeDete_eML_blk_kernel_kernel5<<>> (*gpu_grayImage, *gpu_expanded); cudaMemcpy(gpu_b_b, (void *)&b_b[0], 9ULL, cudaMemcpyHostToDevice); edgeDete_eML_blk_kernel_kernel6<<>> (*gpu_b_b, *gpu_expanded, *gpu_V); edgeDete_eML_blk_kernel_kernel7<<>> (threshold, *gpu_V, *gpu_edgeImage, *gpu_H); cudaMemcpy(&edgeImage[0], gpu_edgeImage, 76800ULL, cudaMemcpyDeviceToHost); cudaFree(*gpu_expanded); cudaFree(*gpu_grayImage); cudaFree(*gpu_b); cudaFree(*gpu_b_b); cudaFree(*gpu_H); cudaFree(*gpu_V); cudaFree(*gpu_edgeImage); } /* Output and update for atomic system: '/Sobel Edge' */ void edgeDetectionModelClass::edgeDetection_SobelEdge(const real_T rtu_grayImage[76800], real_T rtu_threshold, B_SobelEdge_edgeDetection_T *localB) { edgeDetection_eML_blk_kernel_b(rtu_grayImage, rtu_threshold, localB->uv); std::memcpy(&localB->edgeImage[0], &localB->uv[0], 76800U * sizeof(uint8_T)); } void edgeDetectionModelClass::edgeDetection_setupGpuResources(void) { } void edgeDetectionModelClass::edgeDetecti_cleanupGpuResources(void) { } /* Model step function */ void edgeDetectionModelClass::step() { char_T *sErr; void *source_R; /* S-Function (sdspwmmfi2): '/From Multimedia File' */ sErr = GetErrorBuffer(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); source_R = (void *)&edgeDetection_B.FromMultimediaFile[0U]; LibOutputs_FromMMFile(&edgeDetection_DW.FromMultimediaFile_HostLib[0U], GetNullPointer(), GetNullPointer(), source_R, GetNullPointer(), GetNullPointer()); if (*sErr != 0) { rtmSetErrorStatus((&edgeDetection_M), sErr); rtmSetStopRequested((&edgeDetection_M), 1); } /* End of S-Function (sdspwmmfi2): '/From Multimedia File' */ /* MATLAB Function: '/RGB to Gray' */ edgeDetection_RGBtoGray(edgeDetection_B.FromMultimediaFile, &edgeDetection_B.sf_RGBtoGray); /* MATLAB Function: '/Sobel Edge' incorporates: * Constant: '/Threshold' */ edgeDetection_SobelEdge(edgeDetection_B.sf_RGBtoGray.gray, edgeDetection_P.Threshold_Value, &edgeDetection_B.sf_SobelEdge); } /* Model initialize function */ void edgeDetectionModelClass::initialize() { { char_T *sErr; edgeDetection_setupGpuResources(); /* Start for S-Function (sdspwmmfi2): '/From Multimedia File' */ sErr = GetErrorBuffer(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); CreateHostLibrary("frommmfile.dll", &edgeDetection_DW.FromMultimediaFile_HostLib[0U]); createAudioInfo(&edgeDetection_DW.FromMultimediaFile_AudioInfo[0U], 0U, 0U, 0.0, 0, 0, 0, 0, GetNullPointer()); createVideoInfo(&edgeDetection_DW.FromMultimediaFile_VideoInfo[0U], 1U, 15.0, 15.000015000015, "RGB ", 1, 3, 320, 240, 0U, 1, 1, GetNullPointer()); if (*sErr == 0) { LibCreate_FromMMFile(&edgeDetection_DW.FromMultimediaFile_HostLib[0U], 0, (void *) "C:\\MATLAB\\R2020b\\matlab\\toolbox \\images\\imdata\\rhinos.avi", 1, "", "", &edgeDetection_DW.FromMultimediaFile_AudioInfo[0U], &edgeDetection_DW.FromMultimediaFile_VideoInfo[0U], 0U, 1U, 1U, 0U, 0U, 1U, 1.0, 9.2233720368547758E+18); } if (*sErr == 0) { LibStart(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); } if (*sErr != 0) { DestroyHostLibrary(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); if (*sErr != 0) { rtmSetErrorStatus((&edgeDetection_M), sErr); rtmSetStopRequested((&edgeDetection_M), 1); } } /* End of Start for S-Function (sdspwmmfi2): '/From Multimedia File' */ } /* InitializeConditions for S-Function (sdspwmmfi2): '/From Multimedia File' */ LibReset(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); } /* Model terminate function */ void edgeDetectionModelClass::terminate() { char_T *sErr; /* Terminate for S-Function (sdspwmmfi2): '/From Multimedia File' */ sErr = GetErrorBuffer(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); LibTerminate(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); if (*sErr != 0) { rtmSetErrorStatus((&edgeDetection_M), sErr); rtmSetStopRequested((&edgeDetection_M), 1); } LibDestroy(&edgeDetection_DW.FromMultimediaFile_HostLib[0U], 0); DestroyHostLibrary(&edgeDetection_DW.FromMultimediaFile_HostLib[0U]); /* End of Terminate for S-Function (sdspwmmfi2): '/From Multimedia File' */ edgeDetecti_cleanupGpuResources(); } /* Constructor */ edgeDetectionModelClass::edgeDetectionModelClass(): edgeDetection_B() ,edgeDetection_DW() ,edgeDetection_M() { /* Currently there is no constructor body generated.*/ } /* Destructor */ edgeDetectionModelClass::~edgeDetectionModelClass() { /* Currently there is no destructor body generated.*/ } /* Real-Time Model get method */ RT_MODEL_edgeDetection_T * edgeDetectionModelClass::getRTM() { return (&edgeDetection_M); }