diff --git a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp --- a/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp @@ -19,6 +19,8 @@ using namespace mlir; +static constexpr const char kSPIRVModule[] = "__spv__"; + namespace { /// Pattern lowering GPU block/thread size/id to loading SPIR-V invocation /// builtin variables. @@ -285,8 +287,11 @@ return moduleOp.emitRemark("match failure: could not selected memory model " "based on 'spv.target_env'"); - auto spvModule = rewriter.create( - moduleOp.getLoc(), addressingModel, memoryModel.getValue()); + // Add a keyword to the module name to avoid symbolic conflict. + auto spvModuleName = StringRef(kSPIRVModule + moduleOp.getName().str()); + auto spvModule = + rewriter.create(moduleOp.getLoc(), addressingModel, + memoryModel.getValue(), spvModuleName); // Move the region from the module op into the SPIR-V module. Region &spvModuleRegion = spvModule.body(); diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -7,7 +7,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_x() kernel @@ -30,7 +30,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_y() kernel @@ -53,7 +53,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_z() kernel @@ -76,7 +76,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_x() kernel attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} { @@ -100,7 +100,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_y() kernel attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { @@ -121,7 +121,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_z() kernel attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { @@ -142,7 +142,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") gpu.module @kernels { gpu.func @builtin_local_id_x() kernel @@ -165,7 +165,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") gpu.module @kernels { gpu.func @builtin_num_workgroups_x() kernel @@ -182,7 +182,7 @@ // ----- module attributes {gpu.container_module} { - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") gpu.module @kernels { gpu.func @builtin_subgroup_id() kernel @@ -198,7 +198,7 @@ // ----- module attributes {gpu.container_module} { - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") gpu.module @kernels { gpu.func @builtin_num_subgroups() kernel @@ -214,7 +214,7 @@ // ----- module attributes {gpu.container_module} { - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 // CHECK: spv.globalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") gpu.module @kernels { gpu.func @builtin_subgroup_size() kernel diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -21,7 +21,7 @@ return } - // CHECK-LABEL: spv.module Logical GLSL450 + // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { // CHECK-DAG: spv.globalVariable @[[NUMWORKGROUPSVAR:.*]] built_in("NumWorkgroups") : !spv.ptr, Input> // CHECK-DAG: spv.globalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr, Input> diff --git a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir --- a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir @@ -8,7 +8,7 @@ max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> } { gpu.module @kernels { - // CHECK-LABEL: spv.module Physical64 OpenCL + // CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL // CHECK: spv.func // CHECK-SAME: {{%.*}}: f32 // CHECK-NOT: spv.interface_var_abi diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -2,7 +2,7 @@ module attributes {gpu.container_module} { gpu.module @kernels { - // CHECK: spv.module Logical GLSL450 { + // CHECK: spv.module @{{.*}} Logical GLSL450 { // CHECK-LABEL: spv.func @basic_module_structure // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>} // CHECK-SAME: {{%.*}}: !spv.ptr [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>} @@ -28,7 +28,7 @@ module attributes {gpu.container_module} { gpu.module @kernels { - // CHECK: spv.module Logical GLSL450 { + // CHECK: spv.module @{{.*}} Logical GLSL450 { // CHECK-LABEL: spv.func @basic_module_structure_preset_ABI // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32 // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>