/*! [config] name: Test int rotate built-in on CL 1.1 clc_version_min: 10 dimensions: 1 [test] name: rotate int1 kernel_name: test_1_rotate_int global_size: 13 0 0 arg_out: 0 buffer int[13] 1 8 1 2 1 -2147483648 -2147483648 536870912 1 -2147483648 1073741824 -2147483648 2147483647 tolerance 0 arg_in: 1 buffer int[13] 1 1 1 1 1 1 1 1 1 1 1 1 -2 arg_in: 2 buffer int[13] 0 3 32 33 320 31 -1 -3 -32 -33 -34 31 -1 !*/ kernel void test_1_rotate_int(global int* out, global int* in0, global int* in1){ out[get_global_id(0)] = rotate(in0[get_global_id(0)], in1[get_global_id(0)]); } ---- blob: 000 [01001032 2a800c14 40000050 00000000] load.denorm.ls2.s32 t0._y__, u0.zzzz, t0.xxxx, void 001 [00811032 15400c14 40000050 00000000] load.denorm.ls2.s32 t1.x___, u0.yyyy, t0.xxxx, void 002 [0100101b 00001804 40010000 00154008] rotate.s32 t0._y__, t1.xxxx, void, t0.yyyy 003 [00800033 00000c14 40000050 00154008] store.denorm.ls2.s32 mem.x___, u0.xxxx, t0.xxxx, t0.yyyy ---- nir: impl __wrapped_test_1_rotate_int { block b0: // preds: 32x3 %0 = @load_global_invocation_id_zero_base () () 32 %1 = load_const (0x00000000) 32 %2 = @load_uniform (%1 (0x0)) (base=4, range=0, dest_type=invalid) 32 %3 = iadd %0.x, %2 32 %4 = load_const (0x00000004) 32 %5 = imul %3, %4 (0x4) 32 %6 = @load_uniform (%1 (0x0)) (base=1, range=0, dest_type=invalid) 32 %7 = @load_global_etna (%6, %5) (access=none, align_mul=4, align_offset=0) 32 %8 = @load_uniform (%1 (0x0)) (base=2, range=0, dest_type=invalid) 32 %9 = @load_global_etna (%8, %5) (access=none, align_mul=4, align_offset=0) 32 %10 = load_const (0x00000020 = 32) 32 %11 = umod %9, %10 (0x20) 32 %12 = ineg %11 32 %13 = ushr %7, %12 32 %14 = iadd %10 (0x20), %11 32 %15 = ishl %7, %14 32 %16 = ior %13, %15 32 %17 = ine32 %11, %1 (0x0) 32 %18 = b32csel %17, %16, %7 32 %19 = urol %7, %11 32 %20 = uge32 %1 (0x0), %11 32 %21 = b32csel %20, %18, %19 32 %22 = @load_uniform (%1 (0x0)) (base=0, range=0, dest_type=invalid) @store_global_etna (%21, %22, %5) (wrmask=x, access=none, align_mul=4, align_offset=0) // succs: b1 block b1: } ---- # Build options: -cl-std=CL1.2 ; SPIR-V ; Version: 1.0 ; Generator: Khronos LLVM/SPIR-V Translator; 14 ; Bound: 42 ; Schema: 0 OpCapability Addresses OpCapability Linkage OpCapability Kernel OpCapability Int8 %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %34 "test_1_rotate_char" %__spirv_BuiltInGlobalInvocationId %40 = OpString "kernel_arg_type.test_1_rotate_char.char*,char*,char*," %41 = OpString "kernel_arg_type_qual.test_1_rotate_char.,,," OpSource OpenCL_C 102000 OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" OpName %test_1_rotate_char "test_1_rotate_char" OpName %out "out" OpName %in0 "in0" OpName %in1 "in1" OpName %entry "entry" OpName %out_addr "out.addr" OpName %in0_addr "in0.addr" OpName %in1_addr "in1.addr" OpName %call "call" OpName %arrayidx "arrayidx" OpName %call1 "call1" OpName %arrayidx2 "arrayidx2" OpName %call3 "call3" OpName %call4 "call4" OpName %arrayidx5 "arrayidx5" OpName %out_0 "out" OpName %in0_0 "in0" OpName %in1_0 "in1" OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import OpDecorate %__spirv_BuiltInGlobalInvocationId Constant OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId OpDecorate %test_1_rotate_char LinkageAttributes "test_1_rotate_char" Export OpDecorate %out Alignment 1 OpDecorate %in0 Alignment 1 OpDecorate %in1 Alignment 1 OpDecorate %out_addr Alignment 4 OpDecorate %in0_addr Alignment 4 OpDecorate %in1_addr Alignment 4 OpDecorate %out_0 Alignment 1 OpDecorate %in0_0 Alignment 1 OpDecorate %in1_0 Alignment 1 %uint = OpTypeInt 32 0 %uchar = OpTypeInt 8 0 %v3uint = OpTypeVector %uint 3 %_ptr_Input_v3uint = OpTypePointer Input %v3uint %void = OpTypeVoid %_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar %__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input %test_1_rotate_char = OpFunction %void DontInline %9 %out = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in1 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %entry = OpLabel %out_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %in0_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %in1_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function OpStore %out_addr %out Aligned 4 OpStore %in0_addr %in0 Aligned 4 OpStore %in1_addr %in1 Aligned 4 %19 = OpLoad %_ptr_CrossWorkgroup_uchar %in0_addr Aligned 4 %20 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call = OpCompositeExtract %uint %20 0 %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %19 %call %23 = OpLoad %uchar %arrayidx Aligned 1 %24 = OpLoad %_ptr_CrossWorkgroup_uchar %in1_addr Aligned 4 %25 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call1 = OpCompositeExtract %uint %25 0 %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %24 %call1 %28 = OpLoad %uchar %arrayidx2 Aligned 1 %call3 = OpExtInst %uchar %1 rotate %23 %28 %30 = OpLoad %_ptr_CrossWorkgroup_uchar %out_addr Aligned 4 %31 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call4 = OpCompositeExtract %uint %31 0 %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %30 %call4 OpStore %arrayidx5 %call3 Aligned 1 OpReturn OpFunctionEnd %34 = OpFunction %void DontInline %9 %out_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in0_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in1_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %38 = OpLabel %39 = OpFunctionCall %void %test_1_rotate_char %out_0 %in0_0 %in1_0 OpReturn OpFunctionEnd ; SPIR-V ; Version: 1.0 ; Generator: Khronos SPIR-V Tools Linker; 0 ; Bound: 42 ; Schema: 0 OpCapability Addresses OpCapability Kernel OpCapability Int8 %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "test_1_rotate_char" %__spirv_BuiltInGlobalInvocationId %4 = OpString "kernel_arg_type.test_1_rotate_char.char*,char*,char*," %5 = OpString "kernel_arg_type_qual.test_1_rotate_char.,,," OpSource OpenCL_C 102000 OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" OpName %test_1_rotate_char "test_1_rotate_char" OpName %out "out" OpName %in0 "in0" OpName %in1 "in1" OpName %entry "entry" OpName %out_addr "out.addr" OpName %in0_addr "in0.addr" OpName %in1_addr "in1.addr" OpName %call "call" OpName %arrayidx "arrayidx" OpName %call1 "call1" OpName %arrayidx2 "arrayidx2" OpName %call3 "call3" OpName %call4 "call4" OpName %arrayidx5 "arrayidx5" OpName %out_0 "out" OpName %in0_0 "in0" OpName %in1_0 "in1" OpDecorate %__spirv_BuiltInGlobalInvocationId Constant OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId OpDecorate %out Alignment 1 OpDecorate %in0 Alignment 1 OpDecorate %in1 Alignment 1 OpDecorate %out_addr Alignment 4 OpDecorate %in0_addr Alignment 4 OpDecorate %in1_addr Alignment 4 OpDecorate %out_0 Alignment 1 OpDecorate %in0_0 Alignment 1 OpDecorate %in1_0 Alignment 1 %uint = OpTypeInt 32 0 %uchar = OpTypeInt 8 0 %v3uint = OpTypeVector %uint 3 %_ptr_Input_v3uint = OpTypePointer Input %v3uint %void = OpTypeVoid %_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar %30 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_Function__ptr_CrossWorkgroup_uchar = OpTypePointer Function %_ptr_CrossWorkgroup_uchar %__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input %test_1_rotate_char = OpFunction %void DontInline %30 %out = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in1 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %entry = OpLabel %out_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %in0_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function %in1_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uchar Function OpStore %out_addr %out Aligned 4 OpStore %in0_addr %in0 Aligned 4 OpStore %in1_addr %in1 Aligned 4 %32 = OpLoad %_ptr_CrossWorkgroup_uchar %in0_addr Aligned 4 %33 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call = OpCompositeExtract %uint %33 0 %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %32 %call %34 = OpLoad %uchar %arrayidx Aligned 1 %35 = OpLoad %_ptr_CrossWorkgroup_uchar %in1_addr Aligned 4 %36 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call1 = OpCompositeExtract %uint %36 0 %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %35 %call1 %37 = OpLoad %uchar %arrayidx2 Aligned 1 %call3 = OpExtInst %uchar %1 rotate %34 %37 %38 = OpLoad %_ptr_CrossWorkgroup_uchar %out_addr Aligned 4 %39 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 %call4 = OpCompositeExtract %uint %39 0 %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uchar %38 %call4 OpStore %arrayidx5 %call3 Aligned 1 OpReturn OpFunctionEnd %2 = OpFunction %void DontInline %30 %out_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in0_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %in1_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar %40 = OpLabel %41 = OpFunctionCall %void %test_1_rotate_char %out_0 %in0_0 %in1_0 OpReturn OpFunctionEnd Kernels: void test_1_rotate_char(__global char* out, __global char* in0, __global char* in1); cl-program-tester: ../mesa/src/compiler/nir/nir_builder.c:96: nir_builder_alu_instr_finish_and_insert: Assertion `src_bit_size == nir_alu_type_get_type_size(op_info->input_types[i])' failed.