WIP: Emit header files for cpp, cuda and metal targets

[aroidzap]

It is quite cumbersome to work without headers in languages like C++ or CUDA.

I tried to add an experimental CodeGenTargets CPPHeader and CUDAHeader.

Everything marked as *ExternCppDecoration is exported in header. In case of functions and structs the ExternCppDecoration is transitively propagated to its function parameters or struct members.

---

* explicitly prefixed by __extern_cpp in slang, or implicitly generated by [CudaKernel], [DllExport] decorations or others

[csyonghe]

I think the changes here are fine.

We do need to expose these as new target, CppHeader and CudaHeader.

You can keep the isEmittingHeader option in the emitter, just route the header target to the same source Emitter and properly set that option.

This way users can just do
slangc a.slang -o a.h

[jkwak-work]

Wouldn't this conflict with another name that actually uses '_' as a part of its name?
You may want to check how "__extern_cpp" avoids the name mangling.

[aroidzap]

But I would like to keep namespaces somehow, if possible.

[aroidzap]

I changed this to use appendScrubbedName, but it still does basically the same - replaces dot with underscore.

[csyonghe]

This shouldn't be necessary because generateName(inst) will already call appendScrubbedName for you to generate a valid name for the target. Why do we need a separate branch here?

[csyonghe]

I suggest using __extenr_cpp in user code and remove this special case logic.

[aroidzap]

generateName(inst) generates names like _S1, _S2 or something_0, something_1 I wanted to get original names from Slang without any suffixes or changes.

[aroidzap]

Main downside of using __extern_cpp is that it does not support namespaces. Unfortunately I think I would really need to drop this (including auto-transitive __extern_cpp for arguments and struct members), otherwise there is a name mismatch between header and source code - for some reason it worked, but it is not a clear solution. Any ideas how to have __extern_cpp that supports namespaces @csyonghe? Because for example, if I want to have CPU/CUDA code (that switches at runtime if cuda is available), I would have name clashes - and namespaces could solve this for me (if they are different for CPU and CUDA sources).

[csyonghe]

How about this:
we allow __extern_cpp on namespaces as well. And __extern_cpp on a decl means that the decl gets a mangled name that is in the form of: parentNamespaceNameThatHasExternCppModifier_DeclName.

So:

__extern_cpp namespace A
{
         namespace B {
              __extern_cpp struct C;
        }
}
Will have the mangled name `A_C`.

[aroidzap]

Yes this would be nice, thanks.

Unfortunately it probably won't fully solve my problem with linking together CPU and CUDA code.

[aroidzap]

I think the changes here are fine.

We do need to expose these as new target, CppHeader and CudaHeader.

You can keep the isEmittingHeader option in the emitter, just route the header target to the same source Emitter and properly set that option.

This way users can just do slangc a.slang -o a.h

I added hpp,h++,hxx for C++ headers and cuh for CUDA headers, to follow the C++ and CUDA sources extensions already defined in Slang.

[jkwak-work]

You will need to add some tests.

[jkwak-work]

I think this should be "h,hpp,inc".
I have never seen a file with an extension like "h++" or "hxx".

[aroidzap]

I just don't want to confuse C .h, .inc vs C++ headers.
I also think that .c++ and .h++ is really not common, but maybe .cc and .hh can be added.

[csyonghe]

I am fine with treating .h as cpp header because we won't be supporting c as target anytime soon. Treating cpp as the default is likely the desired behavior anyways.

[jkwak-work]

I like to suggest to write some test cases for this part, then it will become more clear of what needs to happen.

[csyonghe]

The PR is showing changes from other recently merged PRs. Make sure you rebase your changes on top of tree, and exclude all the commits that are not part of this change so we can have a clean diff here.

[csyonghe]

I think we should bottleneck name generating logic through the same branch regardless of shouldEmitOnlyHeader setting so our cpp code can actually match the header code. Is there a reason why we need different behavior then just calling generateName here?

[csyonghe]

Instead of skipping here, we should just make sure we never generate the work to emit the struct from the beginning. See CLikeSourceEmitter::computeEmitActions where we are trying to figure out what insts needs to be emitted. We should just make sure the irrelavent structs are not added. Also if you have:

struct MyType {}

[HeaderExport]
void myFunc(MyType t);

That should result MyType to be emitted automatically. By carefully handling computeEmitActions we should be able to handle this gracefully.

[aroidzap]

I have an issue with the #4042 MR and its lowerBuiltinTypesForKernelEntryPoints.

Even though I have something like:

__extern_cpp struct externCppStruct 
{
    __extern_cpp uint2 xy;
};

[CudaKernel]
__extern_cpp void externCppKernel(externCppStruct s) { }

The generated code is now:

struct _VectorStorage_uint2_0
{
    FixedArray<uint, 2>  data_0;
};

struct U_StructStorage_externCppStruct_0
{
    _VectorStorage_uint2_0 xy;
};

__global__ void externCppKernel(U_StructStorage_externCppStruct_0 s_0)
{
    return;
}

Is this expected behavior, even when the struct is marked as __extern_cpp? It is hard to generate useful headers now (the names are not predictable and can change based on other unrelated slang code - especially that number suffix).

Previously it was simply:

struct externCppStruct
{
    uint2  xy;
};

__global__ void externCppKernel(externCppStruct s_0)
{
    return;
}

[csyonghe]

@saipraveenb25 I think the fix here is that if we are lowering a struct type that has __extern_cpp modifier, we want to transfer the modifier to the lowered type.

[csyonghe]

@aroidzap If instead we generate:

struct externCppStruct
{
    FixedArray<uint,2> xy;
};

__global__ void externCppKernel(externCppStruct s_0)
{
    return;
}

will that work for you?

[csyonghe]

The reason here is that we have to lower all types that appear in the parameter list of a CudaKernel to not use the builtin vector types because CUDA builtin vector types has a wider alignment than all other targets, causing mismatch in data layout with C++/HLSL/SPIRV. Our solution is to always lower them into raw arrays.

[aroidzap]

Yes itn would help, thanks.

Regarding to the automatic transitive application of __extern_cpp to function arguments and struct members, do you think i should drop it and just put __extern_cpp everywhere in user code? I would need to apply this also in the source (not only header) but this would require taking back the slangc option or slang attribute.

And regarding the namespaces missing in the __extern_cpp, what do you think is the best idea to approach that?

Thank you very much for your help!

[csyonghe]

I think transitive application to struct members makes sense. I am not sure why we need to apply to argument names, they shouldn't matter, right?

For name mangling, you can check the mangleName function and see how __extern_cpp modifier is handled there.

[aroidzap]

I thought that it would be nice to have __extern_cpp automatically propagated to its argument types, so the user don't have to care at all. But having this just for the structs could be nice a compromise.

What about this namespace support for __extern_cpp? Could it be achieved or I should somehow add prefix by some macro in user code (so it can have different prefix for cuda, cpp and possibly metal code).

[csyonghe]

That looks fine. I don't think there is an existing PR that changes the externCpp behavior for struct fields.

[csyonghe]

This feels hacky. We should use the same name emit logic for header and cpp files. For things that needs to be emitted without name mangling, they should be marked as extern_cpp.

[csyonghe]

We shouldn't use hacks like this. If entry points need to be omitted from header emit, they should be removed from the ir before emitting. We should use a clean up pass before emit to remove entry point functions and all function bodies, then add export decoration to all remaining functions or types so they are not dead code eliminated or ignored by emit.

[csyonghe]

We should use a work list instead of recursive calls to prevent stack overflow.

[csyonghe]

This should be IRArrayTypeBase.

Should also handle IRPtrTypeBase, IRHLSLStructuredBufferTyoeBase etc.

[csyonghe]

Need keep alive and HLSL export.

[csyonghe]

Should delete all function modules after marking everything.