| 654 | } |
| 655 | |
| 656 | void ProgramBagCL::SampleImageD(CLTexImage *dst, CLTexImage *src, int log_scale) |
| 657 | { |
| 658 | cl_kernel kernel; |
| 659 | cl_int w = dst->GetImgWidth(), h = dst->GetImgHeight(); |
| 660 | if(log_scale == 1) |
| 661 | { |
| 662 | kernel = s_sampling->_kernel; |
| 663 | clSetKernelArg(kernel, 0, sizeof(cl_mem), &(src->_clData)); |
| 664 | clSetKernelArg(kernel, 1, sizeof(cl_mem), &(dst->_clData)); |
| 665 | clSetKernelArg(kernel, 2, sizeof(cl_int), &(w)); |
| 666 | clSetKernelArg(kernel, 3, sizeof(cl_int), &(h)); |
| 667 | }else |
| 668 | { |
| 669 | cl_int fullstep = (1 << log_scale); |
| 670 | cl_int halfstep = fullstep >> 1; |
| 671 | kernel = s_sampling_k->_kernel; |
| 672 | clSetKernelArg(kernel, 0, sizeof(cl_mem), &(src->_clData)); |
| 673 | clSetKernelArg(kernel, 1, sizeof(cl_mem), &(dst->_clData)); |
| 674 | clSetKernelArg(kernel, 2, sizeof(cl_int), &(w)); |
| 675 | clSetKernelArg(kernel, 3, sizeof(cl_int), &(h)); |
| 676 | clSetKernelArg(kernel, 4, sizeof(cl_int), &(fullstep)); |
| 677 | clSetKernelArg(kernel, 5, sizeof(cl_int), &(halfstep)); |
| 678 | } |
| 679 | size_t dim0 = 128, dim1 = 1; |
| 680 | //while( w * h / dim0 / dim1 < 8 && dim1 > 1) dim1 /= 2; |
| 681 | size_t gsz[2] = {(w + dim0 - 1) / dim0 * dim0, (h + dim1 - 1) / dim1 * dim1}, lsz[2] = {dim0, dim1}; |
| 682 | cl_int status = clEnqueueNDRangeKernel(_queue, kernel, 2, NULL, gsz, lsz, 0, NULL, NULL); |
| 683 | CheckErrorCL(status, "ProgramBagCL::SampleImageD"); |
| 684 | } |
| 685 | |
| 686 | void ProgramBagCL::FilterInitialImage(CLTexImage* tex, CLTexImage* buf) |
| 687 | { |
no test coverage detected