Ссылка .ll файлы, созданные путем компиляции .cu файла с помощью clang - PullRequest
Новая
       47

Ссылка .ll файлы, созданные путем компиляции .cu файла с помощью clang

0 голосов
/ 11 ноября 2019

Я компилирую следующий код, используя clang с: clang++ -std=c++11 -emit-llvm -c -S $1 --cuda-gpu-arch=sm_30. Это создает файлы vectoradd-cuda-nvptx64-nvidia-cuda-sm_30.ll и vectoradd.ll. Цель запуска некоторого анализа LLVM - ядро, которое, возможно, будет его использовать. Поэтому я хотел бы связать ИК пост-анализа с исполняемым файлом, но я не уверен, как. Когда я пытаюсь связать .ll файлы с llvm-link, я получаю сообщение об ошибке Linking globals named '_Z9vectoraddPiS_S_i': symbol multiply defined!. Я не совсем уверен, как этого добиться, поэтому любая помощь приветствуется. 

#define THREADS_PER_BLOCK 512

__global__ void vectoradd(int *A, int *B, int *C, int N) {
  int gi = threadIdx.x + blockIdx.x * blockDim.x;
  if ( gi < N) {
    C[gi] = A[gi] + B[gi];
  }
}

int main(int argc, char **argv) {
  int N = 10000, *d_A, *d_B, *d_C;

  /// allocate host memory
  std::vector<int> A(N);
  std::vector<int> B(N);
  std::vector<int> C(N);

  /// allocate device memory
  cudaMalloc((void **) &d_A, N * sizeof(int));
  cudaMalloc((void **) &d_B, N * sizeof(int));
  cudaMalloc((void **) &d_C, N * sizeof(int));

  /// populate host data
  for ( size_t i = 0; i < N; ++i) {
    A[i] = i; B[i] = i;
  }

  /// copy to device
  cudaMemcpy(d_A, &A[0], N * sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, &B[0], N * sizeof(int), cudaMemcpyHostToDevice);

  dim3 block(THREADS_PER_BLOCK, 1, 1);
  dim3 grid((N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK, 1, 1);

  vectoradd<<<grid,block>>>(d_A, d_B, d_C, N);
  cudaDeviceSynchronize();

  cudaMemcpy(&C[0], d_C, N * sizeof(int), cudaMemcpyDeviceToHost);

  return 0;
}

1 Ответ

1 голос
/ 11 ноября 2019

Траектория компиляции CUDA в Clang довольно сложная (как в цепочке инструментов NVIDIA), и то, что вы пытаетесь сделать, не может работать. IR LLVM из каждой ветви процесса компиляции должен оставаться отдельным, пока не станут доступны объекты с прямой связью. В результате есть много промежуточных шагов, которые вам нужно будет выполнить вручную.

IR-код LLVM для GPU должен быть сначала скомпилирован в код PTX, а затем собран в двоичную полезную нагрузку, которая может быть связана с хостом. объектные файлы.

Итак, в вашем примере вы сначала делаете что-то вроде: 

clang++ -std=c++11 -emit-llvm -c -S test.cu --cuda-gpu-arch=sm_52

, которое испускает два ИК-файла llvm test-cuda-nvptx64-nvidia-cuda-sm_52.ll и test.ll. Затем код GPU необходимо скомпилировать в PTX (подробнее о бэкэнде nvptx здесь ): 

llc -mcpu=sm_52 test-cuda-nvptx64-nvidia-cuda-sm_52.ll -o test.ptx

Теперь код PTX можно собрать в файл ELF, который позже можно будетсвязаны с помощью nvcc (или компоновщика хоста с парой дополнительных шагов) обычным способом: 

ptxas --gpu-name=sm_52 test.ptx -o test.ptx.o
fatbinary --cuda -64 --create test.fatbin --image=profile=sm_52,file=test.ptx.o

Для кода хоста вы делаете что-то вроде 

llc test.ll
clang -m64 -c test.s

для получения вывода на ассемблереиз IR LLVM и затем соберите его в объектный файл.

Теперь с помощью файла Fatbin, содержащего CUDA скомпилированный код, и объектного файла, содержащего скомпилированный код хоста, вы можете выполнить связывание. Я не смог протестировать связывание объектного файла хоста с Fatbinary с помощью Clang, это то, что вам нужно будет решить самостоятельно. Будет полезно изучить как подробный вывод clang во время вызова компиляции CUDA, так и документацию nvcc, чтобы лучше понять, как работает система построения кода устройства.

...