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: if (arg >= 1 && arg1 == 0) { 13; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_arg3, MemRef_arg3, (arg) * sizeof(double), cudaMemcpyHostToDevice)); 14; CODE-NEXT: { 15; CODE-NEXT: dim3 k0_dimBlock(32); 16; CODE-NEXT: dim3 k0_dimGrid(arg >= 1048545 ? 32768 : (arg + 31) / 32); 17; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_arg3, dev_MemRef_arg2, arg, arg1); 18; CODE-NEXT: cudaCheckKernel(); 19; CODE-NEXT: } 20 21; CODE: cudaCheckReturn(cudaMemcpy(MemRef_arg2, dev_MemRef_arg2, (arg) * sizeof(double), cudaMemcpyDeviceToHost)); 22; CODE-NEXT cudaCheckReturn(cudaFree(dev_MemRef_arg3)); 23; CODE-NEXT cudaCheckReturn(cudaFree(dev_MemRef_arg2)); 24 25; CODE: # kernel0 26; CODE-NEXT: for (int c0 = 0; c0 <= (arg - 32 * b0 - 1) / 1048576; c0 += 1) 27; CODE-NEXT: if (arg >= 32 * b0 + t0 + 1048576 * c0 + 1) 28; CODE-NEXT: Stmt_bb6(0, 32 * b0 + t0 + 1048576 * c0); 29 30; IR-LABEL: call i8* @polly_initContextCUDA() 31; IR: sext i32 %arg to i64 32; IR-NEXT: mul i64 33; IR-NEXT: @polly_allocateMemoryForDevice 34 35target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" 36target triple = "x86_64-unknown-linux-gnu" 37 38define void @hoge(i32 %arg, i32 %arg1, [1000 x double]* %arg2, double* %arg3) { 39bb: 40 br label %bb4 41 42bb4: ; preds = %bb13, %bb 43 br label %bb6 44 45bb5: ; preds = %bb13 46 ret void 47 48bb6: ; preds = %bb6, %bb4 49 %tmp = phi i64 [ 0, %bb4 ], [ %tmp10, %bb6 ] 50 %tmp7 = getelementptr inbounds double, double* %arg3, i64 %tmp 51 %tmp8 = load double, double* %tmp7, align 8 52 %tmp9 = getelementptr inbounds [1000 x double], [1000 x double]* %arg2, i64 0, i64 %tmp 53 store double %tmp8, double* %tmp9, align 8 54 %tmp10 = add nuw nsw i64 %tmp, 1 55 %tmp11 = zext i32 %arg to i64 56 %tmp12 = icmp ne i64 %tmp10, %tmp11 57 br i1 %tmp12, label %bb6, label %bb13 58 59bb13: ; preds = %bb6 60 %tmp14 = zext i32 %arg1 to i64 61 %tmp15 = icmp ne i64 0, %tmp14 62 br i1 %tmp15, label %bb4, label %bb5 63} 64