diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -476,6 +476,11 @@
   }];
   let constructor = "mlir::createConvertGPUToSPIRVPass()";
   let dependentDialects = ["spirv::SPIRVDialect"];
+  let options = [
+    Option<"use64bitIndex", "use-64bit-index",
+           "bool", /*default=*/"false",
+           "Use 64-bit integers to convert index types">
+  ];
 }
 
 //===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -85,7 +85,9 @@
     std::unique_ptr<ConversionTarget> target =
         SPIRVConversionTarget::get(targetAttr);
 
-    SPIRVTypeConverter typeConverter(targetAttr);
+    SPIRVConversionOptions options;
+    options.use64bitIndex = this->use64bitIndex;
+    SPIRVTypeConverter typeConverter(targetAttr, options);
     typeConverter.addConversion([&](gpu::MMAMatrixType type) -> Type {
       return convertMMAToSPIRVType(type);
     });
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
@@ -1,6 +1,9 @@
-// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
+// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=false" %s -o - | FileCheck %s --check-prefix=INDEX32
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_x
@@ -8,14 +11,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_id x
       gpu.return
     }
@@ -24,7 +27,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     %c256 = arith.constant 256 : i32
@@ -34,14 +40,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.block_id y
       gpu.return
     }
@@ -50,7 +56,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_z
@@ -58,14 +67,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.block_id z
       gpu.return
     }
@@ -74,7 +83,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
@@ -82,7 +94,7 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {
@@ -90,7 +102,7 @@
       // Note that this ignores the workgroup size specification in gpu.launch.
       // We may want to define gpu.workgroup_size and convert it to the entry
       // point ABI we want here.
-      // CHECK: spirv.Constant 32 : i32
+      // INDEX32: spirv.Constant 32 : i32
       %0 = gpu.block_dim x
       gpu.return
     }
@@ -99,7 +111,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
@@ -107,12 +122,12 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // The constant value is obtained from the spirv.entry_point_abi.
-      // CHECK: spirv.Constant 4 : i32
+      // INDEX32: spirv.Constant 4 : i32
       %0 = gpu.block_dim y
       gpu.return
     }
@@ -121,7 +136,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
@@ -129,12 +147,12 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
       // The constant value is obtained from the spirv.entry_point_abi.
-      // CHECK: spirv.Constant 1 : i32
+      // INDEX32: spirv.Constant 1 : i32
       %0 = gpu.block_dim z
       gpu.return
     }
@@ -143,7 +161,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_local_id_x
@@ -151,14 +172,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_local_id_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[LOCALINVOCATIONID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.thread_id x
       gpu.return
     }
@@ -167,7 +188,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_num_workgroups_x
@@ -175,14 +199,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMWORKGROUPS]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.grid_dim x
       gpu.return
     }
@@ -191,14 +215,17 @@
 
 // -----
 
-module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") : !spirv.ptr<i32, Input>
   gpu.module @kernels {
     gpu.func @builtin_subgroup_id() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
-      // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPID]]
+      // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_id : index
       gpu.return
     }
@@ -207,14 +234,17 @@
 
 // -----
 
-module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") : !spirv.ptr<i32, Input>
   gpu.module @kernels {
     gpu.func @builtin_num_subgroups() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
-      // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[NUMSUBGROUPS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.num_subgroups : index
       gpu.return
     }
@@ -223,7 +253,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
@@ -231,14 +264,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}}
-  // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+  // INDEX32-LABEL:  spirv.module @{{.*}}
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_dim x
       gpu.return
     }
@@ -247,7 +280,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
@@ -255,14 +291,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}}
-  // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+  // INDEX32-LABEL:  spirv.module @{{.*}}
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.block_dim y
       gpu.return
     }
@@ -271,7 +307,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
@@ -279,14 +318,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}}
-  // CHECK: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+  // INDEX32-LABEL:  spirv.module @{{.*}}
+  // INDEX32: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[WORKGROUPSIZE]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.block_dim z
       gpu.return
     }
@@ -295,7 +334,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_x
@@ -303,14 +345,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_global_id_x() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.global_id x
       gpu.return
     }
@@ -319,7 +361,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_y
@@ -327,14 +372,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_global_id_y() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.global_id y
       gpu.return
     }
@@ -343,7 +388,10 @@
 
 // -----
 
-module attributes {gpu.container_module} {
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_z
@@ -351,14 +399,14 @@
     return
   }
 
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId")
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
   gpu.module @kernels {
     gpu.func @builtin_global_id_z() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
-      // CHECK-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[GLOBALINVOCATIONID]]
+      // INDEX32-NEXT: [[VEC:%.*]] = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32-NEXT: {{%.*}} = spirv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.global_id z
       gpu.return
     }
@@ -368,14 +416,17 @@
 
 // -----
 
-module attributes {gpu.container_module} {
-  // CHECK-LABEL:  spirv.module @{{.*}} Logical GLSL450
-  // CHECK: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
+module attributes {
+  gpu.container_module,
+  spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader, Int64], []>, #spirv.resource_limits<>>
+} {
+  // INDEX32-LABEL:  spirv.module @{{.*}} Logical GLSL450
+  // INDEX32: spirv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") : !spirv.ptr<i32, Input>
   gpu.module @kernels {
     gpu.func @builtin_subgroup_size() kernel
       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
-      // CHECK: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
-      // CHECK-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
+      // INDEX32: [[ADDRESS:%.*]] = spirv.mlir.addressof [[SUBGROUPSIZE]]
+      // INDEX32-NEXT: {{%.*}} = spirv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_size : index
       gpu.return
     }