[CUDA] Upstream the SPIR-V translator changes into KhronosGroup/SPIRV-LLVM-Translator

[AlexeySachkov]

#1091 changed some things in the translator, in particular 09b3b2e.
All these commits were squashed into 7a9a425

We might lose this changes while doing pulldowns from KhronosGroup/SPIRV-LLVM-Translator repo

[AlexeySachkov]

I've found a bug in this commit, which affects translation from SPIR-V to LLVM IR: some built-in SPIR-V variables translated not into OpenCL built-ins, but into something strange, like: __spirv_BuiltInWorkgroupId -> __spirv_WorkgroupId_z (not get_group_id).

Faulty code is additional lines in the following mapping:

llvm/llvm-spirv/lib/SPIRV/OCLUtil.h

Lines 594 to 644 in 3b6799a

	/// Map OpenCL work functions to SPIR-V builtin variables.
	template <>
	inline void SPIRVMap<std::string, SPIRVBuiltinVariableKind>::init() {
	add("get_work_dim", BuiltInWorkDim);
	add("__spirv_GlobalSize_x", BuiltInGlobalSize);
	add("__spirv_GlobalSize_y", BuiltInGlobalSize);
	add("__spirv_GlobalSize_z", BuiltInGlobalSize);
	add("get_global_size", BuiltInGlobalSize);
	add("__spirv_GlobalInvocationId_x", BuiltInGlobalInvocationId);
	add("__spirv_GlobalInvocationId_y", BuiltInGlobalInvocationId);
	add("__spirv_GlobalInvocationId_z", BuiltInGlobalInvocationId);
	add("get_global_id", BuiltInGlobalInvocationId);
	add("__spirv_GlobalOffset_x", BuiltInGlobalOffset);
	add("__spirv_GlobalOffset_y", BuiltInGlobalOffset);
	add("__spirv_GlobalOffset_z", BuiltInGlobalOffset);
	add("get_global_offset", BuiltInGlobalOffset);
	add("__spirv_WorkgroupSize_x", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_y", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_z", BuiltInWorkgroupSize);
	add("get_local_size", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_x", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_y", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_z", BuiltInWorkgroupSize);
	add("get_enqueued_local_size", BuiltInEnqueuedWorkgroupSize);
	add("__spirv_LocalInvocationId_x", BuiltInLocalInvocationId);
	add("__spirv_LocalInvocationId_y", BuiltInLocalInvocationId);
	add("__spirv_LocalInvocationId_z", BuiltInLocalInvocationId);
	add("get_local_id", BuiltInLocalInvocationId);
	add("__spirv_NumWorkgroups_x", BuiltInNumWorkgroups);
	add("__spirv_NumWorkgroups_y", BuiltInNumWorkgroups);
	add("__spirv_NumWorkgroups_z", BuiltInNumWorkgroups);
	add("get_num_groups", BuiltInNumWorkgroups);
	add("__spirv_WorkgroupId_x", BuiltInWorkgroupId);
	add("__spirv_WorkgroupId_y", BuiltInWorkgroupId);
	add("__spirv_WorkgroupId_z", BuiltInWorkgroupId);
	add("get_group_id", BuiltInWorkgroupId);
	add("__spirv_WorkgroupId_x", BuiltInWorkgroupId);
	add("__spirv_WorkgroupId_y", BuiltInWorkgroupId);
	add("__spirv_WorkgroupId_z", BuiltInWorkgroupId);
	add("get_global_linear_id", BuiltInGlobalLinearId);
	add("get_local_linear_id", BuiltInLocalInvocationIndex);
	add("__spirv_LocalInvocationId_x", BuiltInLocalInvocationId);
	add("__spirv_LocalInvocationId_y", BuiltInLocalInvocationId);
	add("__spirv_LocalInvocationId_z", BuiltInLocalInvocationId);
	add("get_sub_group_size", BuiltInSubgroupSize);
	add("get_max_sub_group_size", BuiltInSubgroupMaxSize);
	add("get_num_sub_groups", BuiltInNumSubgroups);
	add("get_enqueued_num_sub_groups", BuiltInNumEnqueuedSubgroups);
	add("get_sub_group_id", BuiltInSubgroupId);
	add("get_sub_group_local_id", BuiltInSubgroupLocalInvocationId);
	}

This map is referenced as:

llvm/llvm-spirv/lib/SPIRV/OCLUtil.h

Lines 130 to 131 in 3b6799a

	typedef SPIRVMap<std::string, SPIRVBuiltinVariableKind>

And the only use of it is in SPIRVReader: built-in id is mapped back to unmangled OpenCL built-in:

llvm/llvm-spirv/lib/SPIRV/SPIRVReader.cpp

Lines 241 to 254 in 3b6799a

	// Variable like GlobalInvolcationId[x] -> get_global_id(x).
	// Variable like WorkDim -> get_work_dim().
	bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV,
	SPIRVBuiltinVariableKind Kind) {
	std::string FuncName = SPIRSPIRVBuiltinVariableMap::rmap(Kind);
	std::string MangledName;
	Type *ReturnTy = GV->getType()->getPointerElementType();
	bool IsVec = ReturnTy->isVectorTy();
	if (IsVec)
	ReturnTy = cast<VectorType>(ReturnTy)->getElementType();
	std::vector<Type *> ArgTy;
	if (IsVec)
	ArgTy.push_back(Type::getInt32Ty(*Context));
	mangleOpenClBuiltin(FuncName, ArgTy, MangledName);

I think it would be better to refactor the whole map, i.e. swap types of key and value.

@Alexander-Johnston, @Ruyk, please fix this bug during upstreaming to KhronosGroup/SPIRV-LLVM-Translator

[AlexeySachkov]

Unfortunately, my analysis wasn't quite correct: mentioned map is actually used to map from string to built-in id:

llvm/llvm-spirv/lib/SPIRV/OCL20ToSPIRV.cpp

Lines 1262 to 1263 in 3b6799a

	if (!SPIRSPIRVBuiltinVariableMap::find(DemangledName.str(), &BVKind))
	continue;

So, it seems that we cannot just remove additional lines with stuff like add("__spirv_GlobalSize_x", BuiltInGlobalSize); - we probably need two independent mappings here:

• LLVM IR -> SPIR-V: it will contain both OpenCL C built-ins as well as ones in form of SPIR-V friendly IR
• SPIR-V -> LLVM IR: it will contain only OpenCL C equivalents

[bader]

+@Naghasan

[AlexeySachkov]

Also, as quick workaround/hack we just need to modify the mapping:

llvm/llvm-spirv/lib/SPIRV/OCLUtil.h

Lines 610 to 616 in 3b6799a

	add("__spirv_WorkgroupSize_x", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_y", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_z", BuiltInWorkgroupSize);
	add("get_local_size", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_x", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_y", BuiltInWorkgroupSize);
	add("__spirv_WorkgroupSize_z", BuiltInWorkgroupSize);

By some reason __spirv_WorkgroupSize_* are inserted twice: if we left only the first one bunch of them, then reverse mapping would only contain BuiltInWorkgroupSize -> "get_local_size", which is fine for the consumption part. The same applies for __spirv_WorkgroupId_* which is also duplicated in this map

[Naghasan]

Oh yes, that's not really the intention here. I will make sure this doesn't happen for SPIR-V -> LLVM IR

[bader]

@AlexeySachkov, @Naghasan, is #1393 resolved this issue?

[Naghasan]

Yes, normally #1393 restored the llvm-spirv to upstream

[bader]

@AlexeySachkov, I'm closing this issue.