[clang] 0a35f44 - [HIP] Perform implicit pointer cast when compiling HIP, not when -fcuda-is-device (#165387)

[via cfe-commits]

Author: Juan Manuel Martinez Caamaño
Date: 2025-11-26T14:02:58+01:00
New Revision: 0a35f44f58f322dece584265e252e21b3ca03530

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

LOG: [HIP] Perform implicit pointer cast when compiling HIP, not when -fcuda-is-device (#165387)

When compiling HIP device code, we add implicit casts for the pointer arguments passed to built-in calls.

When compiling for the host, apply the same casts, since the device side of the source (device functions and kernels) should still pass type checks.

Added: 
    

Modified: 
    clang/lib/Sema/SemaExpr.cpp
    clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index c808dec12a6cf..cfabd1b76c103 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6736,14 +6736,13 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
 
     checkDirectCallValidity(*this, Fn, FD, ArgExprs);
 
-    // If this expression is a call to a builtin function in HIP device
-    // compilation, allow a pointer-type argument to default address space to be
-    // passed as a pointer-type parameter to a non-default address space.
-    // If Arg is declared in the default address space and Param is declared
-    // in a non-default address space, perform an implicit address space cast to
-    // the parameter type.
-    if (getLangOpts().HIP && getLangOpts().CUDAIsDevice && FD &&
-        FD->getBuiltinID()) {
+    // If this expression is a call to a builtin function in HIP compilation,
+    // allow a pointer-type argument to default address space to be passed as a
+    // pointer-type parameter to a non-default address space. If Arg is declared
+    // in the default address space and Param is declared in a non-default
+    // address space, perform an implicit address space cast to the parameter
+    // type.
+    if (getLangOpts().HIP && FD && FD->getBuiltinID()) {
       for (unsigned Idx = 0; Idx < ArgExprs.size() && Idx < FD->param_size();
           ++Idx) {
         ParmVarDecl *Param = FD->getParamDecl(Idx);

diff  --git a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
index 366278f648939..b49c1866caa1c 100644
--- a/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
+++ b/clang/test/SemaHIP/amdgpu-gfx950-load-to-lds.hip
@@ -1,7 +1,7 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx950 -verify=device %s -fcuda-is-device
-// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify=host %s
-// device-no-diagnostics
+// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx950 -verify %s -fcuda-is-device
+// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s
+// expected-no-diagnostics
 
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
@@ -20,11 +20,11 @@ __device__ void i_am_device(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ v
     __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0);
     __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0);
 
-    __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
+    __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0);
 
     __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0);
     __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0);
@@ -46,11 +46,11 @@ __global__ void i_am_kernel(void* src, __amdgpu_buffer_rsrc_t rsrc, __shared__ v
     __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 12, vindex, voffset, soffset, 0, 0);
     __builtin_amdgcn_struct_ptr_buffer_load_lds(rsrc, dst, 16, vindex, voffset, soffset, 0, 0);
 
-    __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
-    __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0); // host-error{{cannot initialize a parameter of type '__attribute__((address_space(3))) void *' with an lvalue of type 'void *'}}
+    __builtin_amdgcn_load_to_lds(src, dst, 1, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 2, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 4, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 12, 0, 0);
+    __builtin_amdgcn_load_to_lds(src, dst, 16, 0, 0);
 
     __builtin_amdgcn_global_load_lds(src, dst, 1, 0 , 0);
     __builtin_amdgcn_global_load_lds(src, dst, 2, 0 , 0);