Kevin Sala Penades
f15904ec71
Until now, strict behavior in the number of threads and blocks has been applied only when the kernel is in bare mode. When this mode is enabled, the values passed in UserNumBlocks and UserThreadLimit are not adjusted and are the definitive values used to launch the kernel. This commit detaches the strictness from the kernel mode. This is going to be used by the kernel replay tool. Additionally, it starts clearing the path for the upcoming OpenMP dims modifier, used to configure multidimensional teams and leagues, which will include strictness choices for teams and threads. All the bare kernels must indicate strict behavior. Asserts are added to check this condition.
75 lines
2.5 KiB
C++
75 lines
2.5 KiB
C++
//===------ API.cpp - Kernel Language (CUDA/HIP) entry points ----- C++ -*-===//
|
|
//
|
|
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
|
// See https://llvm.org/LICENSE.txt for license information.
|
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
//
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
#include "Shared/APITypes.h"
|
|
|
|
#include <cstdio>
|
|
|
|
struct dim3 {
|
|
unsigned x = 0, y = 0, z = 0;
|
|
};
|
|
|
|
struct __omp_kernel_t {
|
|
dim3 __grid_size;
|
|
dim3 __block_size;
|
|
size_t __shared_memory;
|
|
|
|
void *__stream;
|
|
};
|
|
|
|
static __omp_kernel_t __current_kernel = {};
|
|
#pragma omp threadprivate(__current_kernel);
|
|
|
|
extern "C" {
|
|
|
|
// TODO: There is little reason we need to keep these names or the way calls are
|
|
// issued. For now we do to avoid modifying Clang's CUDA codegen. Unclear when
|
|
// we actually need to push/pop configurations.
|
|
unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
|
|
size_t __shared_memory, void *__stream) {
|
|
__omp_kernel_t &__kernel = __current_kernel;
|
|
__kernel.__grid_size = __grid_size;
|
|
__kernel.__block_size = __block_size;
|
|
__kernel.__shared_memory = __shared_memory;
|
|
__kernel.__stream = __stream;
|
|
return 0;
|
|
}
|
|
|
|
unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
|
|
size_t *__shared_memory, void *__stream) {
|
|
__omp_kernel_t &__kernel = __current_kernel;
|
|
*__grid_size = __kernel.__grid_size;
|
|
*__block_size = __kernel.__block_size;
|
|
*__shared_memory = __kernel.__shared_memory;
|
|
*((void **)__stream) = __kernel.__stream;
|
|
return 0;
|
|
}
|
|
|
|
int __tgt_target_kernel(void *Loc, int64_t DeviceId, int32_t NumTeams,
|
|
int32_t ThreadLimit, const void *HostPtr,
|
|
KernelArgsTy *Args);
|
|
|
|
unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
|
|
void *args, size_t sharedMem, void *stream) {
|
|
KernelArgsTy Args = {};
|
|
Args.DynCGroupMem = sharedMem;
|
|
Args.UserNumBlocks[0] = gridDim.x;
|
|
Args.UserNumBlocks[1] = gridDim.y;
|
|
Args.UserNumBlocks[2] = gridDim.z;
|
|
Args.UserThreadLimit[0] = blockDim.x;
|
|
Args.UserThreadLimit[1] = blockDim.y;
|
|
Args.UserThreadLimit[2] = blockDim.z;
|
|
Args.ArgPtrs = reinterpret_cast<void **>(args);
|
|
Args.Flags.IsCUDA = true;
|
|
Args.Flags.StrictBlocksAndThreads = true;
|
|
return __tgt_target_kernel(nullptr, 0, gridDim.x, blockDim.x, func, &Args);
|
|
}
|
|
}
|