cpp-taskflow/taskflow/cuda/algorithm/transform.hpp at master · ezhangle/cpp-taskflow · GitHub

History

282 lines (222 loc) · 7.13 KB

Raw

100

101

102

103

104

105

106

107

108

109

110

111

112

113

114

115

116

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

169

170

171

172

173

174

175

176

177

178

179

180

181

182

183

184

185

186

187

188

189

190

191

192

193

194

195

196

197

198

199

200

201

202

203

204

205

206

207

208

209

210

211

212

213

214

215

216

217

218

219

220

221

222

223

224

225

226

227

228

229

230

231

232

233

234

235

236

237

238

239

240

241

242

243

244

245

246

247

248

249

250

251

252

253

254

255

256

257

258

259

260

261

262

263

264

265

266

267

268

269

270

271

272

273

274

275

276

#pragma once

#include "../cudaflow.hpp"

/**

@file taskflow/cuda/algorithm/transform.hpp

@brief cuda parallel-transform algorithms include file

namespace tf {

// ----------------------------------------------------------------------------

// transform

// ----------------------------------------------------------------------------

namespace detail {

/**

@private

template <size_t nt, size_t vt, typename I, typename O, typename C>

__global__ void cuda_transform_kernel(I first, unsigned count, O output, C op) {

auto tid = threadIdx.x;

auto bid = blockIdx.x;

auto tile = cuda_get_tile(bid, nt*vt, count);

cuda_strided_iterate<nt, vt>(

[=]__device__(auto, auto j) {

auto offset = j + tile.begin;

*(output + offset) = op(*(first+offset));

tid,

tile.count()

);

}

/**

@private

template <size_t nt, size_t vt, typename I1, typename I2, typename O, typename C>

__global__ void cuda_transform_kernel(

I1 first1, I2 first2, unsigned count, O output, C op

) {

auto tid = threadIdx.x;

auto bid = blockIdx.x;

auto tile = cuda_get_tile(bid, nt*vt, count);

cuda_strided_iterate<nt, vt>(

[=]__device__(auto, auto j) {

auto offset = j + tile.begin;

*(output + offset) = op(*(first1+offset), *(first2+offset));

tid,

tile.count()

);

}

} // end of namespace detail -------------------------------------------------

// ----------------------------------------------------------------------------

// CUDA standard algorithms: transform

// ----------------------------------------------------------------------------

/**

@brief performs asynchronous parallel transforms over a range of items

@tparam P execution policy type

@tparam I input iterator type

@tparam O output iterator type

@tparam C unary operator type

@param p execution policy

@param first iterator to the beginning of the range

@param last iterator to the end of the range

@param output iterator to the beginning of the output range

@param op unary operator to apply to transform each item

This method is equivalent to the parallel execution of the following loop on a GPU:

@code{.cpp}

while (first != last) {

*output++ = op(*first++);

}

@endcode

template <typename P, typename I, typename O, typename C>

void cuda_transform(P&& p, I first, I last, O output, C op) {

using E = std::decay_t<P>;

unsigned count = std::distance(first, last);

if(count == 0) {

return;

}

detail::cuda_transform_kernel<E::nt, E::vt, I, O, C>

<<<E::num_blocks(count), E::nt, 0, p.stream()>>> (

first, count, output, op

);

}

/**

@brief performs asynchronous parallel transforms over two ranges of items

@tparam P execution policy type

@tparam I1 first input iterator type

@tparam I2 second input iterator type

@tparam O output iterator type

@tparam C binary operator type

@param p execution policy

@param first1 iterator to the beginning of the first range

@param last1 iterator to the end of the first range

@param first2 iterator to the beginning of the second range

@param output iterator to the beginning of the output range

@param op binary operator to apply to transform each pair of items

This method is equivalent to the parallel execution of the following loop on a GPU:

@code{.cpp}

while (first1 != last1) {

*output++ = op(*first1++, *first2++);

}

@endcode

template <typename P, typename I1, typename I2, typename O, typename C>

void cuda_transform(

P&& p, I1 first1, I1 last1, I2 first2, O output, C op

) {

using E = std::decay_t<P>;

unsigned count = std::distance(first1, last1);

if(count == 0) {

return;

}

detail::cuda_transform_kernel<E::nt, E::vt, I1, I2, O, C>

<<<E::num_blocks(count), E::nt, 0, p.stream()>>> (

first1, first2, count, output, op

);

}

// ----------------------------------------------------------------------------

// cudaFlow

// ----------------------------------------------------------------------------

// Function: transform

template <typename I, typename O, typename C>

cudaTask cudaFlow::transform(I first, I last, O output, C c) {

using E = cudaDefaultExecutionPolicy;

unsigned count = std::distance(first, last);

// TODO:

//if(count == 0) {

// return;

//}

return kernel(

E::num_blocks(count), E::nt, 0,

detail::cuda_transform_kernel<E::nt, E::vt, I, O, C>,

first, count, output, c

);

}

// Function: transform

template <typename I1, typename I2, typename O, typename C>

cudaTask cudaFlow::transform(I1 first1, I1 last1, I2 first2, O output, C c) {

using E = cudaDefaultExecutionPolicy;

unsigned count = std::distance(first1, last1);

// TODO:

//if(count == 0) {

// return;

//}

return kernel(

E::num_blocks(count), E::nt, 0,

detail::cuda_transform_kernel<E::nt, E::vt, I1, I2, O, C>,

first1, first2, count, output, c

);

}

// Function: update transform

template <typename I, typename O, typename C>

void cudaFlow::transform(cudaTask task, I first, I last, O output, C c) {

using E = cudaDefaultExecutionPolicy;

unsigned count = std::distance(first, last);

// TODO:

//if(count == 0) {

// return;

//}

kernel(task,

E::num_blocks(count), E::nt, 0,

detail::cuda_transform_kernel<E::nt, E::vt, I, O, C>,

first, count, output, c

);

}

// Function: update transform

template <typename I1, typename I2, typename O, typename C>

void cudaFlow::transform(

cudaTask task, I1 first1, I1 last1, I2 first2, O output, C c

) {

using E = cudaDefaultExecutionPolicy;

unsigned count = std::distance(first1, last1);

// TODO:

//if(count == 0) {

// return;

//}

kernel(task,

E::num_blocks(count), E::nt, 0,

detail::cuda_transform_kernel<E::nt, E::vt, I1, I2, O, C>,

first1, first2, count, output, c

);

}

// ----------------------------------------------------------------------------

// cudaFlowCapturer

// ----------------------------------------------------------------------------

// Function: transform

template <typename I, typename O, typename C>

cudaTask cudaFlowCapturer::transform(I first, I last, O output, C op) {

return on([=](cudaStream_t stream) mutable {

cudaDefaultExecutionPolicy p(stream);

cuda_transform(p, first, last, output, op);

});

}

// Function: transform

template <typename I1, typename I2, typename O, typename C>

cudaTask cudaFlowCapturer::transform(

I1 first1, I1 last1, I2 first2, O output, C op

) {

return on([=](cudaStream_t stream) mutable {

cudaDefaultExecutionPolicy p(stream);

cuda_transform(p, first1, last1, first2, output, op);

});

}

// Function: transform

template <typename I, typename O, typename C>

void cudaFlowCapturer::transform(

cudaTask task, I first, I last, O output, C op

) {

on(task, [=] (cudaStream_t stream) mutable {

cudaDefaultExecutionPolicy p(stream);

cuda_transform(p, first, last, output, op);

});

}

// Function: transform

template <typename I1, typename I2, typename O, typename C>

void cudaFlowCapturer::transform(

cudaTask task, I1 first1, I1 last1, I2 first2, O output, C op

) {

on(task, [=] (cudaStream_t stream) mutable {

cudaDefaultExecutionPolicy p(stream);

cuda_transform(p, first1, last1, first2, output, op);

});

}

} // end of namespace tf -----------------------------------------------------

cpp-taskflow/taskflow/cuda/algorithm/transform.hpp at master · ezhangle/cpp-taskflow · GitHub

Expand file tree

Breadcrumbs

transform.hpp

Latest commit

History

Breadcrumbs

transform.hpp

File metadata and controls