test/GPGPU/size-cast.ll

   1 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
   2 ; RUN: -disable-output < %s | \
   3 ; RUN: FileCheck -check-prefix=CODE %s
   4
   5 ; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
   6 ; RUN: FileCheck %s -check-prefix=IR
   7
   8 ; REQUIRES: pollyacc
   9
  10 ; This test case ensures that we properly sign-extend the types we are using.
  11
  12 ; CODE: Code
  13 ; CODE-NEXT: ====
  14 ; CODE-NEXT: # host
  15 ; CODE-NEXT: if (arg >= 1 && arg1 == 0) {
  16 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_arg3, MemRef_arg3, (arg) * sizeof(double), cudaMemcpyHostToDevice));
  17 ; CODE-NEXT:   {
  18 ; CODE-NEXT:     dim3 k0_dimBlock(32);
  19 ; CODE-NEXT:     dim3 k0_dimGrid(arg >= 1048546 ? 32768 : floord(arg + 31, 32));
  20 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_arg3, dev_MemRef_arg2, arg, arg1);
  21 ; CODE-NEXT:     cudaCheckKernel();
  22 ; CODE-NEXT:   }
  23
  24 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_arg2, dev_MemRef_arg2, (arg) * sizeof(double), cudaMemcpyDeviceToHost));
  25 ; CODE-NEXT: }
  26
  27 ; CODE: # kernel0
  28 ; CODE-NEXT: for (int c0 = 0; c0 <= (arg - 32 * b0 - 1) / 1048576; c0 += 1)
  29 ; CODE-NEXT:   if (arg >= 32 * b0 + t0 + 1048576 * c0 + 1)
  30 ; CODE-NEXT:     Stmt_bb6(0, 32 * b0 + t0 + 1048576 * c0);
  31
  32 ; IR:        call i8* @polly_initContextCUDA()
  33 ; IR-NEXT:   sext i32 %arg to i64
  34 ; IR-NEXT:   mul i64
  35 ; IR-NEXT:   @polly_allocateMemoryForDevice
  36
  37 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
  38 target triple = "x86_64-unknown-linux-gnu"
  39
  40 define void @hoge(i32 %arg, i32 %arg1, [1000 x double]* %arg2, double* %arg3) {
  41 bb:
  42   br label %bb4
  43
  44 bb4:                                              ; preds = %bb13, %bb
  45   br label %bb6
  46
  47 bb5:                                              ; preds = %bb13
  48   ret void
  49
  50 bb6:                                              ; preds = %bb6, %bb4
  51   %tmp = phi i64 [ 0, %bb4 ], [ %tmp10, %bb6 ]
  52   %tmp7 = getelementptr inbounds double, double* %arg3, i64 %tmp
  53   %tmp8 = load double, double* %tmp7, align 8
  54   %tmp9 = getelementptr inbounds [1000 x double], [1000 x double]* %arg2, i64 0, i64 %tmp
  55   store double undef, double* %tmp9, align 8
  56   %tmp10 = add nuw nsw i64 %tmp, 1
  57   %tmp11 = zext i32 %arg to i64
  58   %tmp12 = icmp ne i64 %tmp10, %tmp11
  59   br i1 %tmp12, label %bb6, label %bb13
  60
  61 bb13:                                             ; preds = %bb6
  62   %tmp14 = zext i32 %arg1 to i64
  63   %tmp15 = icmp ne i64 0, %tmp14
  64   br i1 %tmp15, label %bb4, label %bb5
  65 }