[HIPIFY] compile texture_driver

[ywj55555]

Excuse me, how to compile HIPIFY/tests/unit_tests/samples/2_Cookbook
/11_texture_driver/ this test. Mainly how to compile and generate tex2dKernel.code. Thanks

[emankov]

Are you asking about compiling the CUDA source file or the hipified HIP source file?

[ywj55555]

Hello, file HIPIFY/tests/unit_tests/samples/2_Cookbook/11_texture_driver
/tex2dKernel.cpp, this is the cuda source code from HIP to CUDA. It can be compiled using nvcc, but there is an extern variable in this file: extern texture<float, 2, cudaReadModeElementType> tex; I want to know how to use NVCC or hipcc to compile and generate tex2dKernel .code, thinks.

[ywj55555]

Hello, file HIPIFY/tests/unit_tests/samples/2_Cookbook/11_texture_driver /tex2dKernel.cpp, this is the cuda source code from HIP to CUDA. It can be compiled using nvcc, but there is an extern variable in this file: extern texture<float, 2, cudaReadModeElementType> tex; I want to know how to use NVCC or hipcc to compile and generate tex2dKernel .code, thinks.

I changed "extern texture<float, 2, cudaReadModeElementType> tex; "in tex2dKernel.cpp to "#if CUDA_VERSION < 12000 texture<float, 2, cudaReadModeElementType> tex;#endif" which is in texture2dDrv.cpp, then delete the definition of the variable tex in texture2dDrv.cpp, and then convert it into a HIP program. Using the hipcc --genco tex2dKernel.cpp -o tex2dKernel .code command to compile successfully, but if using extern variables in tex2dKernel.cpp, I don’t know how to generate tex2dKernel .code. How can I achieve this? Thanks.

[ywj55555]

The tex variable is not used in texture2dDrv.cpp. texture2dDrv.cpp only obtains tex in the form of a string through cuModuleGetTexRef(&texref, Module, "tex"); therefore, should the declaration and definition of tex be placed in tex2dKernel. cpp file? Thanks.

[emankov]

It is unclear, why you want to compile this particular hipify-clang unit test, which is dedicated to internal hipification testing only, and doesn't even have an entry point.

Anyway, you may do it by the following additional steps:

1. Add an entry point to the tex2dKernel.cpp, for instance:

int main(int argc, char** argv) {
  float out = 0.f;
  tex2dKernel<<<512, 512>>>(&out, 128, 128);
  return 0;
}

2. Compile the tex2dKernel.cpp by nvcc < 12.0.0 (because of the deleted texture support in CUDA 12.0):

   a) without additional options to obtain executable;
   b) with --cubin to generate a Cubin from PTX intermediate file;
   c) with --fatbin to generate a Fat binary from the source.

or:

2. Hipify the changed tex2dKernel.cpp source by hipify-clang.
3. Compile the hipified tex2dKernel.cpp.hip either by hipcc script or directly by clang.