| 45 | |
| 46 | template<typename T> |
| 47 | void nonMaximal(float* x_out, float* y_out, float* resp_out, unsigned* count, |
| 48 | const unsigned idim0, const unsigned idim1, const T* resp_in, |
| 49 | const unsigned edge, const unsigned max_corners) { |
| 50 | auto nonMax = |
| 51 | common::getKernel("arrayfire::cuda::nonMax", {{susan_cuh_src}}, |
| 52 | TemplateArgs(TemplateTypename<T>())); |
| 53 | |
| 54 | dim3 threads(BLOCK_X, BLOCK_Y); |
| 55 | dim3 blocks(divup(idim0 - edge * 2, BLOCK_X), |
| 56 | divup(idim1 - edge * 2, BLOCK_Y)); |
| 57 | |
| 58 | auto d_corners_found = memAlloc<unsigned>(1); |
| 59 | CUDA_CHECK(cudaMemsetAsync(d_corners_found.get(), 0, sizeof(unsigned), |
| 60 | getActiveStream())); |
| 61 | |
| 62 | EnqueueArgs qArgs(blocks, threads, getActiveStream()); |
| 63 | |
| 64 | nonMax(qArgs, x_out, y_out, resp_out, d_corners_found.get(), idim0, idim1, |
| 65 | resp_in, edge, max_corners); |
| 66 | POST_LAUNCH_CHECK(); |
| 67 | |
| 68 | CUDA_CHECK(cudaMemcpyAsync(count, d_corners_found.get(), sizeof(unsigned), |
| 69 | cudaMemcpyDeviceToHost, getActiveStream())); |
| 70 | CUDA_CHECK(cudaStreamSynchronize(cuda::getActiveStream())); |
| 71 | } |
| 72 | |
| 73 | } // namespace kernel |
| 74 | } // namespace cuda |
nothing calls this directly
no test coverage detected