| 208 | } |
| 209 | |
| 210 | result |
| 211 | launch_kernel_from_so(omp_sscp_executable_object::omp_sscp_kernel *kernel, |
| 212 | const rt::range<3> &num_groups, |
| 213 | const rt::range<3> &local_size, unsigned shared_memory, |
| 214 | void **kernel_args) { |
| 215 | // *** Do NOT change these values without changing also on the compiler side |
| 216 | // in host/StaticLocalMemoryPass.cpp *** |
| 217 | // for internal use in group algorithms |
| 218 | constexpr std::size_t internal_local_memory = 1024 * sizeof(uint64_t); |
| 219 | // for local memory global variables |
| 220 | constexpr std::size_t static_local_mem_size = 1024 * 64 * sizeof(uint64_t); |
| 221 | std::size_t total_internal_local_mem_size = |
| 222 | internal_local_memory + static_local_mem_size; |
| 223 | |
| 224 | if (num_groups.size() == 1 && shared_memory == 0) { |
| 225 | // still need to be able to support group algorithms |
| 226 | // make thread-local in case we have multiple threads submitting. |
| 227 | // |
| 228 | // Note: This data array is also utilized to implement |
| 229 | // static local memory (i.e. globals in address space 3 of fixed size). |
| 230 | // At offset 0, internal memory is used for group algorithms, |
| 231 | // starting at offset 1024*sizeof(uint64_t) it is |
| 232 | // used for such static local memory. |
| 233 | static thread_local std::vector<char> internal_local_memory; |
| 234 | auto aligned_internal_local_memory = resize_and_strongly_align( |
| 235 | internal_local_memory, total_internal_local_mem_size); |
| 236 | |
| 237 | omp_sscp_executable_object::work_group_info info{ |
| 238 | num_groups, rt::id<3>{0, 0, 0}, local_size, nullptr, |
| 239 | aligned_internal_local_memory}; |
| 240 | kernel(&info, kernel_args); |
| 241 | return make_success(); |
| 242 | } |
| 243 | |
| 244 | #ifndef _OPENMP |
| 245 | HIPSYCL_DEBUG_WARNING << "omp_queue: SSCP kernel launching was built without OpenMP " |
| 246 | "support, the kernel will execute sequentially!" |
| 247 | << std::endl; |
| 248 | #endif |
| 249 | |
| 250 | #ifdef _OPENMP |
| 251 | #pragma omp parallel |
| 252 | #endif |
| 253 | { |
| 254 | // get page aligned local memory from heap |
| 255 | static thread_local std::vector<char> local_memory; |
| 256 | static thread_local std::vector<char> internal_local_memory; |
| 257 | auto aligned_local_memory = |
| 258 | resize_and_strongly_align(local_memory, shared_memory); |
| 259 | auto aligned_internal_local_memory = resize_and_strongly_align( |
| 260 | internal_local_memory, total_internal_local_mem_size); |
| 261 | #ifdef _OPENMP |
| 262 | #pragma omp for collapse(3) |
| 263 | #endif |
| 264 | for (std::size_t k = 0; k < num_groups.get(2); ++k) { |
| 265 | for (std::size_t j = 0; j < num_groups.get(1); ++j) { |
| 266 | for (std::size_t i = 0; i < num_groups.get(0); ++i) { |
| 267 | omp_sscp_executable_object::work_group_info info{ |
no test coverage detected