# \[NVPTX\] llc -march=nvptx64 -mcpu=sm\_20 generates invalid zero align for device function params

**URL:** <https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778>\
**Category:** LLVM Dev List Archives\
**Created:** [July 11, 2012, 1:46am UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778 "2012-07-11T01:46:44Z")\
**Posts on this page:** 10\
**Page:** 1

<div class="post-metadata">

**Author:** ![Dmitry\_N\_Mikushin](https://avatars.discourse-cdn.com/v4/letter/d/e5b9ba/32.png) [@Dmitry\_N\_Mikushin](https://discourse.llvm.org/u/Dmitry_N_Mikushin)\
**Post date:** [July 11, 2012, 1:46am UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/1 "2012-07-11T01:46:44Z")

</div>

Hello,

FYI, this is a bug [http://llvm.org/bugs/show\_bug.cgi?id=13324](http://llvm.org/bugs/show_bug.cgi?id=13324)

When compiling the following code for sm\_20, func params are by some reason given with .align 0, which is invalid. Problem does not occur if compiled for sm\_10.

> cat test.ll  
> ; ModuleID = ‘\_\_kernelgen\_main\_module’  
> target datalayout = “e-p:64:64-i64:64:64-f64:64:64-n1:8:16:32:64”  
> target triple = “ptx64-unknown-unknown”

%struct.float2 = type { float, float }

define ptx\_device void @\_\_internal\_dsmul(%struct.float2\* noalias nocapture sret %agg.result, %struct.float2\* nocapture byval %x, %struct.float2\* nocapture byval %y) nounwind inlinehint alwaysinline {  
entry:  
%y1 = getelementptr inbounds %struct.float2\* %x, i64 0, i32 1  
%0 = load float\* %y1, align 4  
%sub = fsub float -0.000000e+00, %0  
%1 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %sub, float 4.097000e+03, float %0) nounwind  
%2 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %0, float 4.097000e+03, float %1) nounwind  
%y5 = getelementptr inbounds %struct.float2\* %y, i64 0, i32 1  
%3 = load float\* %y5, align 4  
%sub7 = fsub float -0.000000e+00, %3  
%4 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %sub7, float 4.097000e+03, float %3) nounwind  
%5 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %3, float 4.097000e+03, float %4) nounwind  
%sub12 = fsub float %0, %2  
%sub14 = fsub float %3, %5  
%6 = tail call float @llvm.nvvm.mul.rn.f(float %0, float %3) nounwind  
%sub18 = fsub float -0.000000e+00, %6  
%7 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %2, float %5, float %sub18) nounwind  
%8 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %2, float %sub14, float %7) nounwind  
%9 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %5, float %sub12, float %8) nounwind  
%10 = tail call float asm “mad.f32 $0, $1, $2, $3;”, “=f,f,f,f”(float %sub12, float %sub14, float %9) nounwind  
%x24 = getelementptr inbounds %struct.float2\* %y, i64 0, i32 0  
%11 = load float\* %x24, align 4  
%12 = tail call float @llvm.nvvm.mul.rn.f(float %0, float %11) nounwind  
%x26 = getelementptr inbounds %struct.float2\* %x, i64 0, i32 0  
%13 = load float\* %x26, align 4  
%14 = tail call float @llvm.nvvm.mul.rn.f(float %13, float %3) nounwind  
%add = fadd float %12, %14  
%add29 = fadd float %10, %add  
%15 = tail call float @llvm.nvvm.add.rn.f(float %6, float %add29) nounwind  
%sub32 = fsub float %6, %15  
%16 = tail call float @llvm.nvvm.add.rn.f(float %sub32, float %add29) nounwind  
%agg.result.0 = getelementptr inbounds %struct.float2\* %agg.result, i64 0, i32 0  
store float %16, float\* %agg.result.0, align 8  
%agg.result.1 = getelementptr inbounds %struct.float2\* %agg.result, i64 0, i32 1  
store float %15, float\* %agg.result.1, align 4  
ret void  
}

declare ptx\_device float @llvm.nvvm.add.rn.f(float, float) nounwind readnone

declare ptx\_device float @llvm.nvvm.mul.rn.f(float, float) nounwind readnone

> llc -march=nvptx64 -mcpu=sm\_20 test.ll -o test.ptx  
> cat test.ptx  
> //  
> // Generated by LLVM NVPTX Back-End  
> //

.version 3.0  
.target sm\_20, texmode\_independent  
.address\_size 64

// .globl \_\_internal\_dsmul  
.func \_\_internal\_dsmul(  
.param .b64 \_\_internal\_dsmul\_param\_0,  
.param .align 0 .b8 \_\_internal\_dsmul\_param\_1[8],  
.param .align 0 .b8 \_\_internal\_dsmul\_param\_2[8]  
) // @\_\_internal\_dsmul  
{  
.reg .pred %p\<396\>;  
.reg .s16 %rc\<396\>;  
.reg .s16 %rs\<396\>;  
.reg .s32 %r\<396\>;  
.reg .s64 %rl\<396\>;  
.reg .f32 %f\<396\>;  
.reg .f64 %fl\<396\>;

// BB#0: // %entry  
mov.b64 %rl0, \_\_internal\_dsmul\_param\_1;  
cvta.local.u64 %rl0, %rl0;  
ld.f32 %f0, [%rl0+4];  
neg.f32 %f1, %f0;  
mov.b64 %rl1, \_\_internal\_dsmul\_param\_2;  
mov.f32 %f2, 0f45800800;  
// inline asm  
mad.f32 %f1, %f1, %f2, %f0;  
// inline asm  
// inline asm  
mad.f32 %f3, %f0, %f2, %f1;  
// inline asm  
cvta.local.u64 %rl1, %rl1;  
ld.f32 %f1, [%rl1+4];  
neg.f32 %f4, %f1;  
// inline asm  
mad.f32 %f4, %f4, %f2, %f1;  
// inline asm  
// inline asm  
mad.f32 %f4, %f1, %f2, %f4;  
// inline asm  
sub.f32 %f5, %f0, %f3;  
sub.f32 %f6, %f1, %f4;  
mul.rn.f32 %f2, %f0, %f1;  
neg.f32 %f7, %f2;  
// inline asm  
mad.f32 %f7, %f3, %f4, %f7;  
// inline asm  
// inline asm  
mad.f32 %f3, %f3, %f6, %f7;  
// inline asm  
// inline asm  
mad.f32 %f3, %f4, %f5, %f3;  
// inline asm  
// inline asm  
mad.f32 %f3, %f5, %f6, %f3;  
// inline asm  
ld.f32 %f4, [%rl1];  
mul.rn.f32 %f0, %f0, %f4;  
ld.f32 %f4, [%rl0];  
mul.rn.f32 %f1, %f4, %f1;  
add.f32 %f0, %f0, %f1;  
add.f32 %f1, %f3, %f0;  
add.rn.f32 %f0, %f2, %f1;  
sub.f32 %f2, %f2, %f0;  
add.rn.f32 %f1, %f2, %f1;  
ld.param.u64 %rl0, [\_\_internal\_dsmul\_param\_0];  
st.f32 [%rl0], %f1;  
st.f32 [%rl0+4], %f0;  
ret;  
}

> ptxas -arch=sm\_20 test.ptx -o ptx.cubin  
> ptxas test.ptx, line 13; error : Alignment must be a power of two  
> ptxas test.ptx, line 14; error : Alignment must be a power of two  
> ptxas fatal : Ptx assembly aborted due to errors

---

<div class="post-metadata">

**Author:** ![Dmitry\_N\_Mikushin](https://avatars.discourse-cdn.com/v4/letter/d/e5b9ba/32.png) [@Dmitry\_N\_Mikushin](https://discourse.llvm.org/u/Dmitry_N_Mikushin)\
**Post date:** [November 9, 2012, 12:16pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/2 "2012-11-09T12:16:32Z")

</div>

Dear all,

I'm attaching a patch that should fix the issue mentioned above. It  
simply makes the same check seen in the same file for global  
variables:

&nbsp;&nbsp;emitPTXAddressSpace(PTy-\>getAddressSpace(), O);  
&nbsp;&nbsp;if (GVar-\>getAlignment() == 0)  
&nbsp;&nbsp;&nbsp;&nbsp;O \<\< " .align " \<\< (int) TD-\>getPrefTypeAlignment(ETy);  
&nbsp;&nbsp;else  
&nbsp;&nbsp;&nbsp;&nbsp;O \<\< " .align " \<\< GVar-\>getAlignment();

Could you please review and commit? Do you think it needs a test case?

Thanks,  
- D.

dmikushin@hp2:~/forge/align0\> llc -march=nvptx64 -mcpu=sm\_20 align0.ll -o -  
//  
// Generated by LLVM NVPTX Back-End  
//

.version 3.1  
.target sm\_20  
.address\_size 64

&nbsp;&nbsp;// .globl \_\_internal\_dsmul  
.visible .func \_\_internal\_dsmul(  
&nbsp;&nbsp;.param .b64 \_\_internal\_dsmul\_param\_0,  
&nbsp;&nbsp;.param .align 4 .b8 \_\_internal\_dsmul\_param\_1[8],  
&nbsp;&nbsp;.param .align 4 .b8 \_\_internal\_dsmul\_param\_2[8]  
) // @\_\_internal\_dsmul  
{  
&nbsp;&nbsp;.reg .pred %p\<396\>;  
&nbsp;&nbsp;.reg .s16 %rc\<396\>;  
&nbsp;&nbsp;.reg .s16 %rs\<396\>;  
&nbsp;&nbsp;.reg .s32 %r\<396\>;  
&nbsp;&nbsp;.reg .s64 %rl\<396\>;  
&nbsp;&nbsp;.reg .f32 %f\<396\>;  
&nbsp;&nbsp;.reg .f64 %fl\<396\>;

// BB#0: // %entry  
&nbsp;&nbsp;mov.b64 %rl0, \_\_internal\_dsmul\_param\_1;  
&nbsp;&nbsp;cvta.local.u64 %rl0, %rl0;  
&nbsp;&nbsp;ld.f32 %f0, [%rl0+4];  
&nbsp;&nbsp;neg.f32 %f1, %f0;  
&nbsp;&nbsp;mov.b64 %rl1, \_\_internal\_dsmul\_param\_2;  
&nbsp;&nbsp;mov.f32 %f2, 0f45800800;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f1, %f1, %f2, %f0;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f3, %f0, %f2, %f1;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;cvta.local.u64 %rl1, %rl1;  
&nbsp;&nbsp;ld.f32 %f1, [%rl1+4];  
&nbsp;&nbsp;neg.f32 %f4, %f1;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f4, %f4, %f2, %f1;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f4, %f1, %f2, %f4;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;sub.f32 %f5, %f0, %f3;  
&nbsp;&nbsp;sub.f32 %f6, %f1, %f4;  
&nbsp;&nbsp;mul.rn.f32 %f2, %f0, %f1;  
&nbsp;&nbsp;neg.f32 %f7, %f2;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f7, %f3, %f4, %f7;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f3, %f3, %f6, %f7;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f3, %f4, %f5, %f3;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;mad.f32 %f3, %f5, %f6, %f3;  
&nbsp;&nbsp;// inline asm  
&nbsp;&nbsp;ld.f32 %f4, [%rl1];  
&nbsp;&nbsp;mul.rn.f32 %f0, %f0, %f4;  
&nbsp;&nbsp;ld.f32 %f4, [%rl0];  
&nbsp;&nbsp;mul.rn.f32 %f1, %f4, %f1;  
&nbsp;&nbsp;add.f32 %f0, %f0, %f1;  
&nbsp;&nbsp;add.f32 %f1, %f3, %f0;  
&nbsp;&nbsp;add.rn.f32 %f0, %f2, %f1;  
&nbsp;&nbsp;sub.f32 %f2, %f2, %f0;  
&nbsp;&nbsp;add.rn.f32 %f1, %f2, %f1;  
&nbsp;&nbsp;ld.param.u64 %rl0, [\_\_internal\_dsmul\_param\_0];  
&nbsp;&nbsp;st.f32 [%rl0], %f1;  
&nbsp;&nbsp;st.f32 [%rl0+4], %f0;  
&nbsp;&nbsp;ret;  
}

[align0.patch](https://discourse.llvm.org/uploads/short-url/yWh7aXv15yg2iYKmlFYNww6kw0o.patch) (517 Bytes)

---

<div class="post-metadata">

**Author:** ![Duncan\_Sands](https://avatars.discourse-cdn.com/v4/letter/d/c6cbf5/32.png) [@Duncan\_Sands](https://discourse.llvm.org/u/Duncan_Sands)\
**Post date:** [November 9, 2012, 12:55pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/3 "2012-11-09T12:55:17Z")

</div>

Hi Dmitry,

\> I'm attaching a patch that should fix the issue mentioned above. It

> simply makes the same check seen in the same file for global  
> variables:
> 
> &nbsp;&nbsp;&nbsp;emitPTXAddressSpace(PTy-\>getAddressSpace(), O);  
> &nbsp;&nbsp;&nbsp;if (GVar-\>getAlignment() == 0)  
> &nbsp;&nbsp;&nbsp;&nbsp;&nbsp;O \<\< " .align " \<\< (int) TD-\>getPrefTypeAlignment(ETy);  
> &nbsp;&nbsp;&nbsp;else  
> &nbsp;&nbsp;&nbsp;&nbsp;&nbsp;O \<\< " .align " \<\< GVar-\>getAlignment();

it's not quite the same because your patch uses the ABI alignment, while  
in this snippet it is the preferred alignment (which is usually the same  
as the ABI alignment, but may be bigger).

> Could you please review and commit? Do you think it needs a test case?

Yes, it needs a testcase.

Ciao, Duncan.

---

<div class="post-metadata">

**Author:** ![Dmitry\_N\_Mikushin](https://avatars.discourse-cdn.com/v4/letter/d/e5b9ba/32.png) [@Dmitry\_N\_Mikushin](https://discourse.llvm.org/u/Dmitry_N_Mikushin)\
**Post date:** [November 9, 2012, 1:17pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/4 "2012-11-09T13:17:51Z")

</div>

Hi Duncan,

You're right, global variables use preferred alignment. And - yes,  
preferred alignment in this case is bigger: 8 instead of 4. NVIDIA's  
prop. compiler gives 4. However, since CUDA 5.0 ptx modules are  
linkable with each other, I think alignments for externally visible  
functions and data should all follow ABI rules.

Is there a guide on making tests? I have ~5 pending patches never  
accepted simply because I'm not familiar with LLVM test system :-/

- D.

---

<div class="post-metadata">

**Author:** ![Duncan\_Sands](https://avatars.discourse-cdn.com/v4/letter/d/c6cbf5/32.png) [@Duncan\_Sands](https://discourse.llvm.org/u/Duncan_Sands)\
**Post date:** [November 9, 2012, 1:23pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/5 "2012-11-09T13:23:11Z")

</div>

Hi Dmitry,

> You're right, global variables use preferred alignment. And - yes,  
> preferred alignment in this case is bigger: 8 instead of 4. NVIDIA's  
> prop. compiler gives 4. However, since CUDA 5.0 ptx modules are  
> linkable with each other, I think alignments for externally visible  
> functions and data should all follow ABI rules.

giving it an alignment of 8 does follow ABI rules: everything that is  
8 byte aligned has an address that is a multiple of 4, i.e. it is also  
4 byte aligned.

It would be wrong to assume that external globals have the preferred  
alignment: they can only be assumed to have the ABI alignment. But  
as we are aligning this variable we can give it the alignment we like  
as long as it is at least 4.

In short, I think using the preferred alignment is correct in this  
context.

> Is there a guide on making tests? I have ~5 pending patches never  
> accepted simply because I'm not familiar with LLVM test system :-/

Your test would go in test/CodeGen/PTX/. It should use FileCheck,  
take a look at test/CodeGen/X86/zext-trunc.ll for an example. As  
there are no PTX tests at all for the moment, you need to create a  
lit.local.cfg file in PTX/. Imitate the X86 one.

Ciao, Duncan.

---

<div class="post-metadata">

**Author:** ![jholewinski](https://sea1.discourse-cdn.com/flex021/user_avatar/discourse.llvm.org/jholewinski/32/14280_2.png) [@jholewinski](https://discourse.llvm.org/u/jholewinski)\
**Post date:** [November 9, 2012, 9:49pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/6 "2012-11-09T21:49:35Z")

</div>

Test cases exist under test/CodeGen/NVPTX (name changed in May). Now that I’m back at NVIDIA, I’m going to be running through the bugzilla issues (thanks Dmitry for the reports!). I have practically the exact same patch here in my queue. 🙂

In this case, I would prefer ABI alignment for compatibility with the vendor compiler. It should work either way, but I do need to audit the codebase and tie up any issues here.

---

<div class="post-metadata">

**Author:** ![Dmitry\_N\_Mikushin](https://avatars.discourse-cdn.com/v4/letter/d/e5b9ba/32.png) [@Dmitry\_N\_Mikushin](https://discourse.llvm.org/u/Dmitry_N_Mikushin)\
**Post date:** [November 9, 2012, 9:58pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/7 "2012-11-09T21:58:33Z")

</div>

Hi Justin,

Very glad you're finally back! How's you PhD?

Please see the attached updated patch. Since you think ABI is better,  
I'm only adding a test case.

- D.

[align0.patch](https://discourse.llvm.org/uploads/short-url/iQYFaqQQvMNiCrHFt10m7wcSgI7.patch) (898 Bytes)

---

<div class="post-metadata">

**Author:** ![Duncan\_Sands](https://avatars.discourse-cdn.com/v4/letter/d/c6cbf5/32.png) [@Duncan\_Sands](https://discourse.llvm.org/u/Duncan_Sands)\
**Post date:** [November 10, 2012, 9:37am UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/8 "2012-11-10T09:37:27Z")

</div>

Hi Justin,

> Test cases exist under test/CodeGen/NVPTX (name changed in May).

I've deleted the empty PTX directory.

&nbsp;&nbsp;&nbsp;Now that I'm

> back at NVIDIA, I'm going to be running through the bugzilla issues (thanks  
> Dmitry for the reports!). I have practically the exact same patch here in my  
> queue. 🙂
> 
> In this case, I would prefer ABI alignment for compatibility with the vendor  
> compiler.

I don't really understand this argument. If the vendor compiler is aligning to  
4 (say) then some globals will have address a multiple of 4, some will have  
address a multiple of 8, some will have address a multiple of 16 etc, depending  
on the accidents of just where in memory they happen to be placed. For  
example, if you have two 4 byte globals that follow each other in memory, and  
that are 4 byte aligned, then if the first one has address a multiple of 4 then  
the second will have address a multiple of 8. In short lots of variables will  
be 8 byte aligned by accident. If LLVM gives them all an alignment of 8, what  
does that change? OK, I will now admit that there is an effect if assumptions  
are being made about globals being placed next to each other: if you declare  
two globals A and B immediately after each other in the IR then the LLVM  
semantics doesn't guarantee that they will be laid out one immediately after  
the other in memory. But that's how it happens in practice so maybe people  
are (wrongly) relying on that. Bumping up the alignment to a multiple of 8  
may add extra padding between A and B, causing B to not be at the position that  
such naughty people are expecting.

&nbsp;&nbsp;&nbsp;It should work either way, but I do need to audit the codebase and

> tie up any issues here.

The IR optimizers already bump the alignment of some globals up to the  
preferred alignment, check out enforceKnownAlignment in Local.cpp (it ends  
up being called from instcombine).

Ciao, Duncan.

---

<div class="post-metadata">

**Author:** ![jholewinski](https://sea1.discourse-cdn.com/flex021/user_avatar/discourse.llvm.org/jholewinski/32/14280_2.png) [@jholewinski](https://discourse.llvm.org/u/jholewinski)\
**Post date:** [November 10, 2012, 2:37pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/9 "2012-11-10T14:37:09Z")

</div>

Perhaps “compatibility” is the wrong term to use here. For now, I would like to “match” what the vendor compiler does. I don’t think using preferred alignment would hurt anything in terms of correctness, but I need to go through the entire back-end to see what effects it could have on performance (e.g. adding extra padding increases local memory usage). It could be a complete non-issue for all I know right now.

---

<div class="post-metadata">

**Author:** ![Duncan\_Sands](https://avatars.discourse-cdn.com/v4/letter/d/c6cbf5/32.png) [@Duncan\_Sands](https://discourse.llvm.org/u/Duncan_Sands)\
**Post date:** [November 10, 2012, 3:38pm UTC](https://discourse.llvm.org/t/nvptx-llc-march-nvptx64-mcpu-sm-20-generates-invalid-zero-align-for-device-function-params/23778/10 "2012-11-10T15:38:25Z")

</div>

Hi Justin,

> Perhaps "compatibility" is the wrong term to use here. For now, I would like to  
> "match" what the vendor compiler does. I don't think using preferred alignment  
> would hurt anything in terms of correctness, but I need to go through the entire  
> back-end to see what effects it could have on performance (e.g. adding extra  
> padding increases local memory usage). It could be a complete non-issue for all  
> I know right now.

don't forget that the backend defines what the preferred alignment is. You can  
set it equal to the ABI alignment. If measurements show that a higher alignment  
is better than you can bump the preferred alignment up to a higher value. What  
I'm saying is that globals should always be output using the preferred alignment  
not the ABI alignment: it is the preferred value itself you should be playing  
with.

Ciao, Duncan.