[HIPIFY] Errors when processing the NVIDIA sample "simpleCubemapTexture"

[wuren2020]

Using HIPIFY to process the simpleCubemapTexture example shipped in the CUDA 10.2 toolkit (commands: cd /path/to/NVIDIA_CUDA-10.2_Samples/0_Simple/simpleCubemapTexture; hipify-clang simpleCubemapTexture.cu -- -I/path/to/NVIDIA_CUDA-10.2_Samples/common/inc) produces the following error:

warning: Unknown CUDA version 10.2. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
In file included from <built-in>:1:
In file included from /usr/lib/llvm-11/lib/clang/11.0.0/include/__clang_cuda_runtime_wrapper.h:333:
/usr/local/cuda-10.2/include/texture_indirect_functions.h:218:3: error: use of undeclared identifier '__nv_tex_surf_handler'
  __nv_tex_surf_handler("__itexCubemap", ptr, obj, x, y, z);
  ^
/usr/local/cuda-10.2/include/texture_indirect_functions.h:228:3: note: in instantiation of function template specialization 'texCubemap<float>' requested here
  texCubemap(&ret, texObject, x, y, z);
  ^
/tmp/simpleCubemapTexture.cu-edba05.hip:104:52: note: in instantiation of function template specialization 'texCubemap<float>' requested here
        g_odata[face*width*width + y*width + x] = -texCubemap<float>(tex, cx, cy, cz);
                                                   ^
1 error generated when compiling for host.
Error while processing /tmp/simpleCubemapTexture.cu-edba05.hip.

Any thoughts will be appreciated.

[wuren2020]

Not sure if this is relevant, but also had issues with other samples as follows.

hipify-clang 0_Simple/simpleSurfaceWrite/simpleSurfaceWrite.cu --skip-excluded-preprocessor-conditional-blocks -- -I/path/to/NVIDIA_CUDA-10.2_Samples/common/incgot the following error:

warning: Unknown CUDA version 10.2. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
/tmp/simpleSurfaceWrite.cu-8cd6ee.hip:226:21: warning: CUDA identifier is deprecated.
    checkCudaErrors(cudaMemcpyToArray(cuArray,
                    ^
/tmp/simpleSurfaceWrite.cu-8cd6ee.hip:68:5: error: use of undeclared identifier 'surf2Dwrite'
    surf2Dwrite(gIData[y * width + x],
    ^
In file included from <built-in>:1:
In file included from /usr/lib/llvm-10/lib/clang/10.0.1/include/__clang_cuda_runtime_wrapper.h:324:
/usr/local/cuda-10.2/include/texture_indirect_functions.h:145:4: error: use of undeclared identifier '__nv_tex_surf_handler'
   __nv_tex_surf_handler("__itex2D", ptr, obj, x, y);
   ^
/usr/local/cuda-10.2/include/texture_indirect_functions.h:154:3: note: in instantiation of function template specialization 'tex2D<float>' requested here
  tex2D(&ret, texObject, x, y);
  ^
/tmp/simpleSurfaceWrite.cu-8cd6ee.hip:96:29: note: in instantiation of function template specialization 'tex2D<float>' requested here
    gOData[y * width + x] = tex2D<float>(tex, tu, tv);
                            ^
1 warning and 2 errors generated when compiling for host.
Error while processing /tmp/simpleSurfaceWrite.cu-8cd6ee.hip.

, and hipify-clang 0_Simple/simpleMultiCopy/simpleMultiCopy.cu --skip-excluded-preprocessor-conditional-blocks -- -I/path/to/NVIDIA_CUDA-10.2_Samples/common/inc got

/tmp/simpleMultiCopy.cu-2fb81b.hip:128:7: error: no matching function for call to 'max'
      max((32.0f / (_ConvertSMVer2Cores(deviceProp.major, deviceProp.minor) *
...
...
/usr/local/cuda-10.2/include/crt/math_functions.hpp:1079:31: note: candidate function not viable: call to __device__ function from __host__ function
__MATH_FUNCTIONS_DECL__ float max(float a, float b)

even though MATH_FUNCTIONS_DECL is defined as __host__ __device__.

[emankov]

@wuren2020,
Could you, please, split this issue based on CUDA SDK sample name and the error reported? Thus, we shall see one issue per sample with the corresponding input command line and error log.

P.S.
Although supporting all CUDA SDK samples is not the first-class priority for hipify-clang, I agree, that we don't want to see clang compilation errors without a chance to hipify them in any way.

[zhangxi19920406]

For the origin issue (error output with "error: use of undeclared identifier '__nv_tex_surf_handler'")
I find many cases in CUDA samples:
eg:
0_Simple
cudaNvSci/imageKernels.cu, simpleCubemapTexture/simpleCubemapTexture.cu, simpleLayeredTexture/simpleLayeredTexture.cu, simplePitchLinearTexture/simplePitchLinearTexture.cu, simpleSurfaceWrite/simpleSurfaceWrite.cu, simpleTexture/simpleTexture.cu, simpleTextureDrv/simpleTexture_kernel.cu,
2_Graphics
bindlessTexture/bindlessTexture_kernel.cu, simpleTexture3D/simpleTexture3D_kernel.cu, volumeFiltering/volumeFilter_kernel.cu, volumeRender/volumeRender_kernel.cu
....

Overall there are around 40 files have this issue

[emankov]

The original bug: https://bugs.llvm.org/show_bug.cgi?id=26400, which was finally fixed in clang only in Sept 2021 by https://reviews.llvm.org/D110089.

hipify-clang should be built against the latest trunk LLVM (currently, 14.0.0git).

The command line to test:

hipify-clang --print-stats --cuda-path="c:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.5" "c:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.5/0_Simple/simpleCubemapTexture/simpleCubemapTexture.cu" --skip-excluded-preprocessor-conditional-blocks -I"c:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.5\common\inc" -- -D_ALLOW_COMPILER_AND_STL_VERSION_MISMATCH

Output:

In file included from C:\Users\TT\AppData\Local\Temp\simpleCubemapTexture.cu-aeae30.hip:31:
In file included from c:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.5\common\inc\helper_functions.h:35:
c:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.5\common\inc\helper_image.h:409:14: warning: 'fopen' is deprecated: This function or variable may be unsafe. Consider using fopen_s instead. To disable
      deprecation, use _CRT_SECURE_NO_WARNINGS. See online help for details. [-Wdeprecated-declarations]
  FILE *fh = fopen(filename, "rb");
             ^
C:\Program Files (x86)\Windows Kits\10\Include\10.0.22000.0\ucrt\stdio.h:212:20: note: 'fopen' has been explicitly marked deprecated here
    _Check_return_ _CRT_INSECURE_DEPRECATE(fopen_s)
                   ^
C:\Program Files\Microsoft Visual Studio\2022\Preview\VC\Tools\MSVC\14.31.30818\include\vcruntime.h:335:55: note: expanded from macro '_CRT_INSECURE_DEPRECATE'
        #define _CRT_INSECURE_DEPRECATE(_Replacement) _CRT_DEPRECATE_TEXT(    \
                                                      ^
C:\Program Files\Microsoft Visual Studio\2022\Preview\VC\Tools\MSVC\14.31.30818\include\vcruntime.h:325:47: note: expanded from macro '_CRT_DEPRECATE_TEXT'
#define _CRT_DEPRECATE_TEXT(_Text) __declspec(deprecated(_Text))
                                              ^
1 warning generated when compiling for host.

[HIPIFY] info: file 'c:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.5/0_Simple/simpleCubemapTexture/simpleCubemapTexture.cu' statistics:
  CONVERTED refs count: 40
  UNCONVERTED refs count: 0
  CONVERSION %: 100.0
  REPLACED bytes: 783
  TOTAL bytes: 8690
  CHANGED lines of code: 35
  TOTAL lines of code: 254
  CODE CHANGED (in bytes) %: 9.0
  CODE CHANGED (in lines) %: 13.8
  TIME ELAPSED s: 6.99
[HIPIFY] info: CONVERTED refs by type:
  device: 3
  memory: 11
  texture: 3
  include_cuda_main_header: 1
  type: 10
  numeric_literal: 9
  define: 1
  kernel_launch: 2
[HIPIFY] info: CONVERTED refs by API:
  CUDA RT API: 40
[HIPIFY] info: CONVERTED refs by names:
  cudaAddressModeWrap: 3
  cudaArray: 1
  cudaArrayCubemap: 1
  cudaChannelFormatDesc: 1
  cudaChannelFormatKindFloat: 1
  cudaCreateChannelDesc: 1
  cudaCreateTextureObject: 1
  cudaDestroyTextureObject: 1
  cudaDeviceProp: 1
  cudaDeviceSynchronize: 2
  cudaFilterModeLinear: 1
  cudaFree: 1
  cudaFreeArray: 1
  cudaGetDeviceProperties: 1
  cudaLaunchKernel: 2
  cudaMalloc: 1
  cudaMalloc3DArray: 1
  cudaMemcpy: 1
  cudaMemcpy3D: 1
  cudaMemcpy3DParms: 1
  cudaMemcpyDeviceToHost: 1
  cudaMemcpyHostToDevice: 1
  cudaReadModeElementType: 1
  cudaResourceDesc: 2
  cudaResourceTypeArray: 1
  cudaTextureDesc: 2
  cudaTextureObject_t: 2
  cuda_runtime.h: 1
  make_cudaExtent: 2
  make_cudaPitchedPtr: 1
  make_cudaPos: 2