KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator
Null constant inserted by `memset` handling not added to entry points interface
- Dominant language
- LLVM
- Stars
- 625
- Forks
- 279
- Avg merge
- 3d 5h
- Merged PRs (30d)
- 34
Description
The handling of `memset` can create new variables which never get added to the entry points interface. In the example below the variable `%30` is not handled correctly.
example opencl C file:
```c
__kernel void test(__global uchar* input, __global int* output)
{
*output = (int [5]){ 0,0,0,0,0, }[*input];
}
```
llvm ir:
```llvm
; ModuleID = ''
source_filename = "test.cl"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
target triple = "spir64-unknown-unknown"
@constinit = private global [5 x i32] zeroinitializer, align 4
; Function Attrs: convergent noinline norecurse nounwind optnone
define dso_local spir_kernel void @resample_ref_14493953481400187574_0(ptr addrspace(1) noundef align 1 %0, ptr addrspace(1) noundef align 4 %1) #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
%3 = alloca ptr addrspace(1), align 8
%4 = alloca ptr addrspace(1), align 8
%5 = alloca [5 x i32], align 4
store ptr addrspace(1) %0, ptr %3, align 8
store ptr addrspace(1) %1, ptr %4, align 8
call void @llvm.memset.p0.i64(ptr align 4 %5, i8 0, i64 20, i1 false)
%6 = getelementptr inbounds [5 x i32], ptr %5, i64 0, i64 0
call void @llvm.memcpy.p0.p0.i64(ptr align 4 %5, ptr align 4 @constinit, i64 20, i1 false)
%7 = load ptr addrspace(1), ptr %3, align 8
%8 = load i8, ptr addrspace(1) %7, align 1
%9 = zext i8 %8 to i64
%10 = getelementptr inbounds [5 x i32], ptr %5, i64 0, i64 %9
%11 = load i32, ptr %10, align 4
%12 = load ptr addrspace(1), ptr %4, align 8
store i32 %11, ptr addrspace(1) %12, align 4
ret void
}
; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write)
declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #1
; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #2
attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
attributes #1 = { nocallback nofree nounwind willreturn memory(argmem: write) }
attributes #2 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
!llvm.module.flags = !{!0, !1}
!opencl.ocl.version = !{!2}
!opencl.spir.version = !{!2}
!llvm.ident = !{!3}
!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 7, !"frame-pointer", i32 2}
!2 = !{i32 3, i32 0}
!3 = !{!"clang version 16.0.6 (Fedora 16.0.6-3.fc38)"}
!4 = !{i32 1, i32 1}
!5 = !{!"none", !"none"}
!6 = !{!"uchar*", !"int*"}
!7 = !{!"", !""}
```
spirv:
```spirv
; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 51
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpCapability Int8
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %46 "resample_ref_14493953481400187574_0" %constinit
OpSource OpenCL_C 300000
OpName %constinit "constinit"
OpName %resample_ref_14493953481400187574_0 "resample_ref_14493953481400187574_0"
OpDecorate %constinit Alignment 4
OpDecorate %resample_ref_14493953481400187574_0 LinkageAttributes "resample_ref_14493953481400187574_0" Export
OpDecorate %14 Alignment 1
OpDecorate %15 Alignment 4
OpDecorate %18 Alignment 8
OpDecorate %19 Alignment 8
OpDecorate %21 Alignment 4
OpDecorate %30 Constant
OpDecorate %47 Alignment 1
OpDecorate %48 Alignment 4
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uchar = OpTypeInt 8 0
%ulong_5 = OpConstant %ulong 5
%ulong_20 = OpConstant %ulong 20
%ulong_0 = OpConstant %ulong 0
%_arr_uint_ulong_5 = OpTypeArray %uint %ulong_5
%_ptr_UniformConstant__arr_uint_ulong_5 = OpTypePointer UniformConstant %_arr_uint_ulong_5
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%12 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar
%_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar
%_ptr_Function__arr_uint_ulong_5 = OpTypePointer Function %_arr_uint_ulong_5
%_ptr_Function_uchar = OpTypePointer Function %uchar
%_arr_uchar_ulong_20 = OpTypeArray %uchar %ulong_20
%_ptr_UniformConstant__arr_uchar_ulong_20 = OpTypePointer UniformConstant %_arr_uchar_ulong_20
%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint
%6 = OpConstantNull %_arr_uint_ulong_5
%constinit = OpVariable %_ptr_UniformConstant__arr_uint_ulong_5 UniformConstant %6
%28 = OpConstantNull %_arr_uchar_ulong_20
%30 = OpVariable %_ptr_UniformConstant__arr_uchar_ulong_20 UniformConstant %28
%resample_ref_14493953481400187574_0 = OpFunction %void DontInline %12
%14 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%15 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%16 = OpLabel
%18 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%19 = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function
%21 = OpVariable %_ptr_Function__arr_uint_ulong_5 Function
%22 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_uchar %18
OpStore %22 %14 Aligned 8
%23 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_uchar %19
OpStore %23 %15 Aligned 8
%25 = OpBitcast %_ptr_Function_uchar %21
%32 = OpBitcast %_ptr_UniformConstant_uchar %30
OpCopyMemorySized %25 %32 %ulong_20 Aligned 4
%35 = OpInBoundsPtrAccessChain %_ptr_Function_uint %21 %ulong_0 %ulong_0
OpCopyMemorySized %21 %constinit %ulong_20 Aligned 4
%36 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_uchar %18
%37 = OpLoad %_ptr_CrossWorkgroup_uchar %36 Aligned 8
%38 = OpLoad %uchar %37 Aligned 1
%39 = OpUConvert %ulong %38
%40 = OpInBoundsPtrAccessChain %_ptr_Function_uint %21 %ulong_0 %39
%41 = OpLoad %uint %40 Aligned 4
%44 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_uint %19
%45 = OpLoad %_ptr_CrossWorkgroup_uint %44 Aligned 8
OpStore %45 %41 Aligned 4
OpReturn
OpFunctionEnd
%46 = OpFunction %void DontInline %12
%47 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%48 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
%49 = OpLabel
%50 = OpFunctionCall %void %resample_ref_14493953481400187574_0 %47 %48
OpReturn
OpFunctionEnd
```
Contributor guide
Research direction
Start by reproducing the issue with the OpenCL C example and compare its LLVM IR and SPIR-V output, focusing on the OpEntryPoint interface and the generated %30 variable. Trace the handling of memset-created variables and determine why the variable is emitted without being added to the entry-point interface. Done means the generated entry point includes all required variables for this case.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- c
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 25/100