search for: x86_sse42_crc32_64_64

...; > So you're saying that in this particular case you'd prefer LLVM passes to > know about the range of these PTX intrinsics, rather than Clang adding them > as metadata? > Yep. ValueTracking.cpp already has some iffy target knowledge (someone sneaked a > direct Intrinsic::x86_sse42_crc32_64_64 check in there), but extending it > to other intrinsics in other targets seems like too much... > That's not iffy. That's exactly how it should work, and we should have more of that. There is a major gotcha and that's dealing with the case where the intrinsics don't exist bec...

...prefer LLVM passes to >>> know about the range of these PTX intrinsics, rather than Clang adding them >>> as metadata? >> >> >> Yep. >> >>> ValueTracking.cpp already has some iffy target knowledge (someone sneaked >>> a direct Intrinsic::x86_sse42_crc32_64_64 check in there), but extending it >>> to other intrinsics in other targets seems like too much... >> >> >> That's not iffy. That's exactly how it should work, and we should have >> more of that. There is a major gotcha and that's dealing with the case wh...

...far read-only and accessed using intrinsics. e.g., > %threadIdx.x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x(). > > One possible approach is to have ValueTracking compute the known > bits of these intrinsics as special cases. This approach is already > taken for the x86_sse42_crc32_64_64 intrinsic. However, this > approach may not be elegant because the ranges of these CUDA special > registers depend on the GPU compute capability specified by > -target-cpu. For instance, blockIdx.x is bounded by 65535 in sm_20 > but 2^31-1 in sm_30. Exposing -target-cpu...

...>> >>> them >> >>> as metadata? >> >> >> >> >> >> Yep. >> >> >> >>> ValueTracking.cpp already has some iffy target knowledge (someone >> >>> sneaked >> >>> a direct Intrinsic::x86_sse42_crc32_64_64 check in there), but >> >>> extending it >> >>> to other intrinsics in other targets seems like too much... >> >> >> >> >> >> That's not iffy. That's exactly how it should work, and we should have >> >> more of th...

...variables. These special registers are so far read-only and accessed using intrinsics. e.g., %threadIdx.x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x(). One possible approach is to have ValueTracking compute the known bits of these intrinsics as special cases. This approach is already taken for the x86_sse42_crc32_64_64 intrinsic. However, this approach may not be elegant because the ranges of these CUDA special registers depend on the GPU compute capability specified by -target-cpu. For instance, blockIdx.x is bounded by 65535 in sm_20 but 2^31-1 in sm_30. Exposing -target-cpu to ValueTracking is probably discour...

...x86_sse42_crc64_8 : > + def int_x86_sse42_crc32_64_8 : > Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i8_ty], > [IntrNoMem]>; > - def int_x86_sse42_crc64_64 : GCCBuiltin<"__builtin_ia32_crc32di">, > + def int_x86_sse42_crc32_64_64 : GCCBuiltin<"__builtin_ia32_crc32di">, > Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty], > [IntrNoMem]>; > } > > Modified: llvm/trunk/lib/Analysis/ValueTracking.cpp > URL: http://llvm.org/viewvc/llvm-project/llvm/trun...