CUDA host device code with two code paths

[jpienaar]

Summary

Allow CUDA host device functions with two code paths using CUDA_ARCH to differentiate between code path being compiled.

For example,

__host__ __device__ void host_device_function(void) {
#ifdef __CUDA_ARCH__
  device_only_function();
#else
  host_only_function();
#endif
}

[rnk]

Can you remind me what the CUDA compilation model is currently? My memory was that the clang driver was eventually going to launch two -cc1 actions, one for device and one for host, presumably with different flags. I would expect that lib/Frontend/InitPreprocessor.cpp would define this macro when targeting the device.

If we're doing a single compilation with a fat object approach, we may need to do something weird to get this right. =/

[jpienaar]

I think your memory is correct (at least thats what I thought too). And
yes, the macro would be defined externally when targeting the device. At
that point we can remove this check for macro definition as we'd then be
able to check the flags directly. So I do see this as a temporary solution
which disrupts as little as possible.

Is the concern here with how the test is written? (i.e., the test
explicitly sets this macro which will in future be set by the compiler
itself). In which case it could be changed to

host device void hd1(void) {
#ifdef CUDA_ARCH

hd1d();
hd1h(); // expected-error {{no matching function}}

#else

hd1d(); // expected-error {{no matching function}}
hd1h();

#endif

hd1hd();
hd1g<<<1, 1>>>(); // expected-error {{reference to __global__

function 'hd1g' in host device function}}
}

[pcc]

There is already a flag for this, -fcuda-is-device. We should make that flag a LangOption and use it for this check.

[rnk]

In D6457#7, @pcc wrote:

There is already a flag for this, -fcuda-is-device. We should make that flag a LangOption and use it for this check.

Right, this seems like the correct approach. Is there concern that it would be too disruptive at this stage to define __CUDA_ARCH__?

[jpienaar]

Creating CUDAIsDevice as language options (using fcuda-is-device frontend option). Using this flag instead of definition of CUDA_ARCH to determine whether host/device compilation is occurring.

[jpienaar]

Thanks, that's a good idea. Is this close to what you in mind?

[rnk]

include/clang/Basic/LangOptions.def

160 ↗

(On Diff #16862)

I think you just want regular LANGOPT, given the description of BENIGN_LANGOPT:

// BENIGN_LANGOPT: for options that don't affect the construction of the AST in
//     any way (that is, the value can be different between an implicit module
//     and the user of that module).

lib/Basic/Targets.cpp

1381 ↗

(On Diff #16862)

"... the NVPTX backend." maybe?

lib/Frontend/InitPreprocessor.cpp

873–878 ↗

(On Diff #16862)

I guess this definition is intended to satisfy targeting hypothetical non-NVPTX targets from CUDA. OK.

lib/Sema/SemaCUDA.cpp

91

The predominant style in clang for eliding text from a quotation is to use "text [...] text". I'm not trying to be pedantic, this actually threw me off when I read it. :)

I guess this applies equally to the CUDA reference quotes above.

[jpienaar]

Made recommended changes.

[rnk]

lgtm

[jpienaar]

Thanks (and the comment about text elision was helpful :) ). I don't have commit access, could you assist me in committing this? Thanks

[rnk]

Closed by commit rL223271 (authored by @rnk).

lib/Sema/SemaCUDA.cpp

 #include "clang/Sema/Sema.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Decl.h"
+#include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/SemaDiagnostic.h"
 #include "llvm/ADT/Optional.h"
 #include "llvm/ADT/SmallVector.h"

       (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
     return true;
 
-  if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice)
-    return true;
+  // CUDA B.1.3 "The __device__ and __host__ qualifiers can be used together
+  // however, in which case the function is compiled for both the host and the
+  // device. The __CUDA_ARCH__ macro... can be used to differentiate code paths
+  // between host and device."
+  bool InDeviceMode = getPreprocessor()
+                          .getIdentifierInfo("__CUDA_ARCH__")
+                          ->hasMacroDefinition();
+  if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
+    if ((InDeviceMode && CalleeTarget != CFT_Device) ||
+        (!InDeviceMode && CalleeTarget != CFT_Host))
+      return true;
+  }
 
   return false;
 }

test/SemaCUDA/function-target.cu

 __host__ __device__ void hd1hd(void);
 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
 
-__host__ __device__ void hd1(void) {
-  hd1h(); // expected-error {{no matching function}}
+// __CUDA_ARCH__ is not normally defined and undefined from within a program,
+// but rather originates from the compiler. It is being explicitly manipulated
+// here to test both compilation paths.
+#undef __CUDA_ARCH__
+__host__ __device__ void hd1_hostmode(void) {
+  hd1h(); // no longer an error as assumed to be in host mode
   hd1d(); // expected-error {{no matching function}}
+#ifdef __CUDA_ARCH__
+  hd1d(); // no longer an error as guarded
+#else
+  hd1h(); // no longer an error as guarded
+#endif
   hd1hd();
   hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
 }
+
+#define __CUDA_ARCH__ 250
+__host__ __device__ void hd1_devicemode(void) {
+  hd1h(); // expected-error {{no matching function}}
+  hd1d(); // no longer an error as assumed to be in device mode
+#ifdef __CUDA_ARCH__
+  hd1d(); // no longer an error as guarded
+#else
+  hd1h(); // no longer an error as guarded
+#endif
+  hd1hd();
+}