[clang] df1b2be - [CUDA] Explicitly construct dim3() return values.

Artem Belevich via cfe-commits cfe-commits at lists.llvm.org
Thu May 25 12:41:58 PDT 2023


Author: Artem Belevich
Date: 2023-05-25T12:41:25-07:00
New Revision: df1b2bef0c7cad11681a02e9e2f816b27fb480a6

URL: https://github.com/llvm/llvm-project/commit/df1b2bef0c7cad11681a02e9e2f816b27fb480a6
DIFF: https://github.com/llvm/llvm-project/commit/df1b2bef0c7cad11681a02e9e2f816b27fb480a6.diff

LOG: [CUDA] Explicitly construct dim3() return values.

Fixes CUDA build break caused by 5c082e7e15e38a2eea1f506725ef

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_intrinsics.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index 0e0dd4674997..3c3948863c1d 100644
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -607,27 +607,27 @@ __device__ inline unsigned __clusterDimIsSpecified() {
 }
 
 __device__ inline dim3 __clusterDim() {
-  return {__nvvm_read_ptx_sreg_cluster_nctaid_x(),
-          __nvvm_read_ptx_sreg_cluster_nctaid_y(),
-          __nvvm_read_ptx_sreg_cluster_nctaid_z()};
+  return dim3(__nvvm_read_ptx_sreg_cluster_nctaid_x(),
+              __nvvm_read_ptx_sreg_cluster_nctaid_y(),
+              __nvvm_read_ptx_sreg_cluster_nctaid_z());
 }
 
 __device__ inline dim3 __clusterRelativeBlockIdx() {
-  return {__nvvm_read_ptx_sreg_cluster_ctaid_x(),
-          __nvvm_read_ptx_sreg_cluster_ctaid_y(),
-          __nvvm_read_ptx_sreg_cluster_ctaid_z()};
+  return dim3(__nvvm_read_ptx_sreg_cluster_ctaid_x(),
+              __nvvm_read_ptx_sreg_cluster_ctaid_y(),
+              __nvvm_read_ptx_sreg_cluster_ctaid_z());
 }
 
 __device__ inline dim3 __clusterGridDimInClusters() {
-  return {__nvvm_read_ptx_sreg_nclusterid_x(),
-          __nvvm_read_ptx_sreg_nclusterid_y(),
-          __nvvm_read_ptx_sreg_nclusterid_z()};
+  return dim3(__nvvm_read_ptx_sreg_nclusterid_x(),
+              __nvvm_read_ptx_sreg_nclusterid_y(),
+              __nvvm_read_ptx_sreg_nclusterid_z());
 }
 
 __device__ inline dim3 __clusterIdx() {
-  return {__nvvm_read_ptx_sreg_clusterid_x(),
-          __nvvm_read_ptx_sreg_clusterid_y(),
-          __nvvm_read_ptx_sreg_clusterid_z()};
+  return dim3(__nvvm_read_ptx_sreg_clusterid_x(),
+              __nvvm_read_ptx_sreg_clusterid_y(),
+              __nvvm_read_ptx_sreg_clusterid_z());
 }
 
 __device__ inline unsigned __clusterRelativeBlockRank() {


        


More information about the cfe-commits mailing list