| 48 | } |
| 49 | |
| 50 | inline __device__ float block_reduce_sum(float val) { |
| 51 | __shared__ float data[THREADS_X * THREADS_Y]; |
| 52 | |
| 53 | unsigned idx = threadIdx.x * blockDim.x + threadIdx.y; |
| 54 | |
| 55 | data[idx] = val; |
| 56 | __syncthreads(); |
| 57 | |
| 58 | for (unsigned i = blockDim.y / 2; i > 0; i >>= 1) { |
| 59 | if (threadIdx.y < i) { data[idx] += data[idx + i]; } |
| 60 | |
| 61 | __syncthreads(); |
| 62 | } |
| 63 | |
| 64 | return data[threadIdx.x * blockDim.x]; |
| 65 | } |
| 66 | |
| 67 | template<typename T> |
| 68 | __global__ void keep_features(float* x_out, float* y_out, float* score_out, |
no outgoing calls
no test coverage detected