Skip to content

Instantly share code, notes, and snippets.

View pashu123's full-sized avatar
😇
Working from home

Prashant Kumar pashu123

😇
Working from home
28 followers · 7 following
View GitHub Profile
@pashu123
pashu123 / dispatch_43.spv
Created February 10, 2023 17:23
; SPIR-V
; Version: 1.3
; Generator: Khronos; 22
; Bound: 132
; Schema: 0
OpCapability Shader
OpCapability GroupNonUniformShuffle
OpExtension "SPV_KHR_storage_buffer_storage_class"
%98 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
@pashu123
pashu123 / err_dump.txt
Created February 10, 2023 15:33
/home/prashant/stable.mlir:793:11: error: failed to materialize conversion for result #0 of operation 'hal.interface.constant.load' that remained live after conversion
%10 = linalg.generic {indexing_maps = [#map4, #map4], iterator_types = ["parallel", "parallel"]} ins(%9 : tensor<2x160xf32>) outs(%8 : tensor<2x160xf32>) {
^
/home/prashant/stable.mlir:24:3: note: called from
func.func @forward(%arg0: tensor<1x4x96x96xf32>, %arg1: tensor<1xf32>, %arg2: tensor<2x64x1024xf32>, %arg3: tensor<f32>) -> tensor<1x4x96x96xf32> {
^
/home/prashant/stable.mlir:793:11: note: see current operation: %20 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
%10 = linalg.generic {indexing_maps = [#map4, #map4], iterator_types = ["parallel", "parallel"]} ins(%9 : tensor<2x160xf32>) outs(%8 : tensor<2x160xf32>) {
^
/home/prashant/stable.mlir:793:11: note: see existing live user here: %28 = "spirv.UConvert"(%20) : (i32) -> i64
@pashu123
pashu123 / err_dump.txt
Created February 9, 2023 05:02
/home/prashant/stable.mlir:793:11: error: failed to materialize conversion for result #0 of operation 'hal.interface.constant.load' that remained live after conversion
%10 = linalg.generic {indexing_maps = [#map4, #map4], iterator_types = ["parallel", "parallel"]} ins(%9 : tensor<2x160xf32>) outs(%8 : tensor<2x160xf32>) {
^
/home/prashant/stable.mlir:24:3: note: called from
func.func @forward(%arg0: tensor<1x4x96x96xf32>, %arg1: tensor<1xf32>, %arg2: tensor<2x64x1024xf32>, %arg3: tensor<f32>) -> tensor<1x4x96x96xf32> {
^
/home/prashant/stable.mlir:793:11: note: see current operation: %20 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
%10 = linalg.generic {indexing_maps = [#map4, #map4], iterator_types = ["parallel", "parallel"]} ins(%9 : tensor<2x160xf32>) outs(%8 : tensor<2x160xf32>) {
^
/home/prashant/stable.mlir:793:11: note: see existing live user here: %28 = "spirv.UConvert"(%20) : (i32) -> i64
@pashu123
pashu123 / err_dump.txt
Created February 9, 2023 04:54
/home/prashant/stable.mlir:765:15: error: failed to legalize operation 'arith.sitofp' that was explicitly marked illegal
%3362 = arith.sitofp %3361 : i64 to f32
^
/home/prashant/stable.mlir:24:3: note: called from
func.func @forward(%arg0: tensor<1x4x96x96xf32>, %arg1: tensor<1xf32>, %arg2: tensor<2x64x1024xf32>, %arg3: tensor<f32>) -> tensor<1x4x96x96xf32> {
^
/home/prashant/stable.mlir:765:15: note: see current operation: %42 = "arith.sitofp"(%41) : (i64) -> f32
%3362 = arith.sitofp %3361 : i64 to f32
^
/home/prashant/stable.mlir:788:10: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_nv = []>>}>
@pashu123
pashu123 / module_forward_dispatch_0_vulkan_spirv_fb.mlir
Created February 8, 2023 17:51
module attributes {hal.device.targets = [#hal.device.target<"vulkan", {executable_targets = [#hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, CooperativeMatrixNV], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers, SPV_NV_cooperative_matrix]>, api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], min_subgroup_size = 32, max_subgroup_size
@pashu123
pashu123 / module_forward_dispatch_42_vulkan_spirv_fb.mlir
Created February 8, 2023 15:43
module attributes {hal.device.targets = [#hal.device.target<"vulkan", {executable_targets = [#hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, CooperativeMatrixNV], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers, SPV_NV_cooperative_matrix]>, api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], min_subgroup_size = 32, max_subgroup_size
@pashu123
pashu123 / module_forward_dispatch_43_vulkan_spirv_fb.mlir
Created February 8, 2023 15:43
module attributes {hal.device.targets = [#hal.device.target<"vulkan", {executable_targets = [#hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, CooperativeMatrixNV], [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers, SPV_NV_cooperative_matrix]>, api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 64], min_subgroup_size = 32, max_subgroup_size
@pashu123
pashu123 / vae_cpu.fx
Created January 25, 2023 15:54
graph():
%arg0_1 : [#users=1] = placeholder[target=arg0_1]
%mul : [#users=1] = call_function[target=torch.ops.aten.mul](args = (%arg0_1, 5.489980785067252), kwargs = {})
%_param_constant0 : [#users=1] = get_attr[target=_param_constant0]
%_param_constant1 : [#users=1] = get_attr[target=_param_constant1]
%convolution : [#users=1] = call_function[target=torch.ops.aten.convolution](args = (%mul, %_param_constant0, %_param_constant1, [1, 1], [0, 0], [1, 1], False, [0, 0], 1), kwargs = {})
%_param_constant2 : [#users=1] = get_attr[target=_param_constant2]
%_param_constant3 : [#users=1] = get_attr[target=_param_constant3]
%convolution_1 : [#users=2] = call_function[target=torch.ops.aten.convolution](args = (%convolution, %_param_constant2, %_param_constant3, [1, 1], [1, 1], [1, 1], False, [0, 0], 1), kwargs = {})
%view : [#users=2] = call_function[target=torch.ops.aten.view](args = (%convolution_1, [1, 32, 16, 4096]), kwargs = {})
@pashu123
pashu123 / cuda.fx
Created January 13, 2023 08:09
graph():
%arg0_1 : [#users=1] = placeholder[target=arg0_1]
%arg1_1 : [#users=1] = placeholder[target=arg1_1]
%arg2_1 : [#users=32] = placeholder[target=arg2_1]
%arg3_1 : [#users=1] = placeholder[target=arg3_1]
%cat : [#users=1] = call_function[target=torch.ops.aten.cat](args = ([%arg0_1, %arg0_1],), kwargs = {})
%expand : [#users=1] = call_function[target=torch.ops.aten.expand](args = (%arg1_1, [2]), kwargs = {})
%arange : [#users=1] = call_function[target=torch.ops.aten.arange](args = (0, 160), kwargs = {dtype: torch.float32, device: cuda:0, pin_memory: False})
%mul : [#users=1] = call_function[target=torch.ops.aten.mul](args = (%arange, -9.210340371976184), kwargs = {})
%div : [#users=1] = call_function[target=torch.ops.aten.div](args = (%mul, 160), kwargs = {})
@pashu123
pashu123 / cpu.fx
Created January 13, 2023 08:08
graph():
%arg0_1 : [#users=1] = placeholder[target=arg0_1]
%arg1_1 : [#users=1] = placeholder[target=arg1_1]
%arg2_1 : [#users=32] = placeholder[target=arg2_1]
%arg3_1 : [#users=1] = placeholder[target=arg3_1]
%cat : [#users=1] = call_function[target=torch.ops.aten.cat](args = ([%arg0_1, %arg0_1],), kwargs = {})
%expand : [#users=1] = call_function[target=torch.ops.aten.expand](args = (%arg1_1, [2]), kwargs = {})
%arange : [#users=1] = call_function[target=torch.ops.aten.arange](args = (0, 160), kwargs = {dtype: torch.float16, device: cpu, pin_memory: False})
%mul : [#users=1] = call_function[target=torch.ops.aten.mul](args = (%arange, -9.210340371976184), kwargs = {})
%div : [#users=1] = call_function[target=torch.ops.aten.div](args = (%mul, 160), kwargs = {})