Currently if p_N is defined to 5 then when a kernel is parsed and output all instances of p_N are replaced by 5.
It would be useful to add a parser option so that the #defines are included in the generated source code, and the compilers variables like p_N are left in the generated source code.
This will make the generated code a lot easier to read.
Unfortunately all defines need to be expanded to apply code transformations
#define
values can be found in two places
> head ~/.occa/cache/c8141715ac4e4272/raw_source.cpp
#define block 256
/* The MIT License (MIT)
*
> cat ~/.occa/cache/c8141715ac4e4272/build.json | jq .kernel.props.defines
{
"block": 256
}
I agree that it is useful to do this to extract actual loop bounds and I imagine the parser is designed to do this. However, it seems unlikely that it is impossible to preserve the definitions in the code. Keeping track of how the values are computed would obviate the need to print their numerical values.
It is harder than it seems to maintain the code and transform it
Defines can be anything, even partial code
#define foo 3 +
for (int i = 0; i < foo 5; ++i) {}
Keeping the #define
lines can also cause issues since the compiler will still run the preprocessor
... and if the Define is just a numerical constant... ?
... and why do we need extra logic for that ....?
This is an example of the output kernel
extern "C" __global__ void _occa_ellipticPreconCoarsenHex3D_0(const int Nelements,
const double * __restrict__ R,
const double * __restrict__ qf,
double * __restrict__ qc) {
{
int e = 0 + blockIdx.x;
__shared__ double s_qfff[8][8][8];
__shared__ double s_qcff[2][8][8];
__shared__ double s_qccf[2][2][8];
__shared__ double s_qccc[2][2][2];
__shared__ double s_R[2][8];
{
int k = 0 + threadIdx.z;
{
int j = 0 + threadIdx.y;
{
int i = 0 + threadIdx.x;
const int id = i + j * 8 + k * 8 * 8 + e * 512;
s_qfff[k][j][i] = qf[id];
if ((k == 0) && (j < 2)) {
s_R[j][i] = R[j * 8 + i];
}
}
}
}
The kernels are losing readability on translation. All these numbers were Defines.
Taking a step back, the reason for the feature request was to
make the generated code a lot easier to read
The generated code is meant for the compiler, not for users (similar to preprocessed outputs by the compiler). If there is a parsing issue, specially since the parser is not yet mature, it could help to look at. However, these issues should be sorted out over time.
If the purpose is to make it easier to debug, I agree we should add #line
to match with the original source code
I agree that for debugging this would be a true bonus - nice idea !
I should have been clearer about the reasoning behind this request: one important use case I envision for the new code generation tools in OCCA 1.0 is as a translator from OKL to native threading language for non-OCCA applications. In that case it will be important to preserve readability as the generated CUDA/OpenCL/HIP/OpenMP code will be separated from the OKL code.
That sounds reasonable
There are a few features that will need to be added first
#define
#defines
is too difficult since we would need to track where they are defined and which scopes they touch after transformsparser::getExpression
, substitute constexpr
with their valuesconst
variablesconstexpr
functions#define
lines to the top like usualThanks for contemplating the feature request !
Most helpful comment
That sounds reasonable
There are a few features that will need to be added first
#define
#defines
is too difficult since we would need to track where they are defined and which scopes they touch after transformsparser::getExpression
, substituteconstexpr
with their valuesconst
variablesconstexpr
functions#define
lines to the top like usual